ó
    Eñiû ã                  óâ  • S SK Jr  S SKrS SKrS SKrS SKrS SKrS SKrS SKrS SK	r	S SK
r
S SKrS SKJr  S SKJrJrJr  S SKJr  S SKJrJrJrJrJrJr  S SKrS SKJr  S SKrS SKrS SKJ s  J!r"  S SK#J$r$  S S	K%J&r&J'r'  S S
K(J)r)  S SK*J+r+  S SK,J-r-J.r.J/r/  S SK0J1r1J2r2J3r3  SSK4J5r5J6r6J7r7J8r8  SSK9J:r:  SSK;J<r<J=r=J>r>J r   SSK?J@r@  SSKAJBrBJCrCJDrDJErE  SSKFJGrG  SSKHJIrI  SSKJJKrK  SSKLJMrM  SSKNJOrOJPrPJQrQJRrRJSrS  SSKTJUrUJVrV  SSKWJXrXJYrYJZrZJ[r[  SSK\J]r]  SSK J^r^J_r_J`r`JaraJbrbJcrcJdrdJereJfrfJgrgJhrhJiriJjrjJkrk  SSKlJmrnJoroJprpJqrq  SSKrJsrs  S S!KtJuru  S S"KvJwrwJxrxJyryJzrzJ{r{J|r|J}r}J~r~JrJ€r€J�r�J‚r‚JƒrƒJ„r„J…r…J†r†  S S#K‡JˆrˆJ‰r‰JŠrŠJ‹r‹JŒrŒJ�r�JŽrŽ  S S$K�J�r�J‘r‘J’r’J“r“J”r”  S S%K•J–r–  \(       a,  S S&K—J˜r˜  S S'K™Jšrš  S S(K›Jœrœ  SS)K=J�r�  S S*KvJžrž  S S+KŸJ r   \" S,5      r¡\GRD                  " \£5      r¤\GRJ                  GRM                  \£S-5      r§\GRJ                  GRM                  \£S.5      r¨\GRJ                  GRM                  \£S/5      r©\@" 5       r?S0rªS1 r«ScS2 jr¬ " S3 S45      r­\" S5      SdS5 j5       r® " S6 S75      r¯\GR`                   " S8 S95      5       r±\GR`                   " S: S;5      5       r²\GR`                   " S< S=\²5      5       r³\GR`                   " S> S?\²5      5       r´        SeS@ jrµ " SA SB\�5      r¶\¶" 5       GRn                  r¸SfSC jr¹SfSD jrºSgSE jr»SfSF jr¼ShSG jr½SiSH jr¾ " SI SJ\{5      r¿SjSK jrÀSkSlSL jjrÁ " SM SN\€5      rÂ\ÂGR‡                  SO5         " SP SQ\Â5      rÄ " SR SS5      rÅ\GR`                   " ST SU5      5       rÆ " SV SW5      rÇ\GR`                   " SX SY5      5       rÈ " SZ S[\z\¿\\É\Ê\É\É4   4   4   5      rË\GR`                   " S\ S]5      5       rÌ " S^ S_\�\¿   5      rÍ " S` Sa\Ž5      rÎSmSb jrÏg)né    )ÚannotationsN)Úabstractmethod)ÚCallableÚIterableÚSequence)Ú	lru_cache)ÚAnyÚcastÚOptionalÚTYPE_CHECKINGÚTypeVarÚUnion)Ú
PRECEDENCE)Úget_interface_for_device)ÚidentityÚpreserve_rng_state)Úis_integer_dtype)Ú
OrderedSet)ÚCeilDivÚFloorDivÚModularIndexing)Úget_triton_versionÚhas_triton_packageÚhas_triton_stable_tma_apié   )Úfree_symbol_is_typeÚ
prefix_strÚsymbol_is_typeÚSymT)ÚValueRangesé   )ÚconfigÚirÚmetricsÚutils)ÚAsyncCompile)Ú	code_hashÚget_pathÚPyCodeCacheÚwrite_atomic)Ú'set_kernel_post_grad_provenance_tracing)ÚDefaultHandler)Útriton_heuristics)Úbenchmarker)ÚAutotuneHintÚDevicePropertiesÚReductionHintÚTRITON_MAX_BLOCKÚTRITON_MAX_RSPLIT)Úget_max_y_gridÚnext_power_of_2)ÚBaseSchedulerNodeÚFusedSchedulerNodeÚ	SchedulerÚSchedulerNode)Úget_broadcasted_shape)Úcache_on_selfÚDelayReplaceLineÚget_bounds_index_exprÚget_fused_kernel_nameÚget_kernel_metadataÚis_welford_reductionÚPlaceholderÚprefix_is_reductionÚ	sympy_dotÚsympy_productÚ
sympy_subsÚtriton_typeÚtriton_version_uses_attrs_dictÚupcast_compute_type)Ú_opsÚReductionTypeÚ	StoreModeÚV)Ú"get_kernel_category_by_source_codeé   )ÚBlockPatternMatcher)ÚArgNameÚBackendFeatureÚConstexprArgÚCSEÚCSEVariableÚDeferredLineÚIndentedBufferÚInplacedBufferÚis_buffer_removedÚOpOverridesÚPythonPrinterÚ
RemovedArgÚSizeArgÚ	TensorArgÚWorkspaceArgÚWorkspaceZeroMode)Úconstant_reprÚIterationRangesÚIterationRangesEntryÚIterationRangesRootÚPartialAccumulateÚ
SIMDKernelÚSIMDScheduling)Ú	config_ofÚequal_1_arg_indicesÚnon_constexpr_signatureÚshould_unwrap_unspec_argÚsignature_to_meta)ÚSymbolicCallArg)Ú
ModuleType©ÚDtypePropagationOpsHandler)ÚShapeEnv)ÚIRNode)ÚBlockShapeType)ÚSIMDKernelFeaturesÚ_TÚ
perf_hintsÚscheduleÚfusioné   c                óL   • U S;   nU(       a  SOSnU S;   a  U SU  S3$ U SU  3$ )N)ÚanyÚmaxÚminÚprodÚtriton_helpersÚtl)r{   r|   Ú.Ú2© )Úreduction_typeÚ
use_helperÚmodules      Ú[/home/mande/repo/quber/.venv/lib/python3.13/site-packages/torch/_inductor/codegen/triton.pyÚget_triton_reduction_functionr‡   †   sE   € ØÐ#@Ñ@€JÞ!+Ñ°€FØ˜Ó'Ø�˜˜>Ð*¨!Ð,Ð,à�˜˜>Ð*Ð+Ð+ó    c                óÚ   • [        U [        R                  5      (       d  g[        U [        R                  5      =(       d+    U R                  =(       a    [        U R                  5      S:H  $ )z�"
Is this expression a Sympy Integer or is it an integer sympy Expr
containing no free symbols. The latter case can happen with Identity expr.
Fr   )Ú
isinstanceÚsympyÚExprÚIntegerÚ
is_integerÚlenÚfree_symbols©Úexprs    r†   Úis_sympy_integer_liker“   �   sL   € ô
 �dœEŸJ™J×'Ñ'ØÜ�dœEŸM™MÓ*÷ Ø�‰×7œC × 1Ñ 1Ó2°aÑ7ðrˆ   c                  óJ   • \ rS rSr% Sr0 rS\S'   0 rS\S'   \S
S j5       r	Sr
g	)ÚOpDtypeSupporté›   z“
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_dtypeszdict[str, bool]Úconvert_outputsc                ó¤   • UR                   n[        [        R                  [        R                  /5      U R
                  U'   X R                  U'   g ©N)Ú__name__r   ÚtorchÚfloat32Úfloat64r—   r˜   )ÚclsÚfuncÚconvert_outputÚop_names       r†   Úregister_upcastÚOpDtypeSupport.register_upcast¤   s;   € à—-‘-ˆÜ(2´E·M±MÄ5Ç=Á=Ð3QÓ(Rˆ×Ñ˜WÑ%Ø'5×Ñ˜GÒ$rˆ   r‚   N)r    zCallable[..., str]r¡   ÚboolÚreturnÚNone)r›   Ú
__module__Ú__qualname__Ú__firstlineno__Ú__doc__r—   Ú__annotations__r˜   Úclassmethodr£   Ú__static_attributes__r‚   rˆ   r†   r•   r•   ›   s1   ‡ ñð
 <>ÐÐ8Ó=Ø')€O�_Ó)àó6ó ó6rˆ   r•   c                 óx   • [        5       (       d  gSSKn [        U R                  R                  S5      (       a  gg)zX
import AttrsDescriptor if the triton version is new enough to have this
class defined.
Ú r   NÚAttrsDescriptorz4from triton.compiler.compiler import AttrsDescriptor)r   Útriton.compiler.compilerÚhasattrÚcompiler)Útritons    r†   Úgen_attr_descriptor_importr¶   «   s3   € ô ×ÑØã#ô ˆv�‰×'Ñ'Ð):×;Ñ;ØEàrˆ   c                  ó  • \ rS rSrSr\" \R                  \R                  /5      r	\" \R                  \R                  \R                  /\	Q5      r\ V VVVs0 s H$  nU[        R                  " [         U    S3SSS9_M&     snnnn r\ V VVVs0 s H2  nU[        R                  " [         U   R%                  5        S3SSS9_M4     snnnn r\SS j5       r\SS	 j5       r\SS
 j5       rSrgs  snnnn f s  snnnn f )ÚTritonSymbolsé¾   zM
Stores sympy.Symbol instances and constants associated with triton codegen.
ÚoffsetT©ÚintegerÚnonnegativeÚBLOCK©r¼   Úpositivec           
     ó"  • SnUR                   nU GHê  n[        U[        R                  5      (       a?  [        R
                  R                  R                  UR                     nUR                  nGO}[        U[        R                  [        R                  [        R                  [        R                  [        R                  [        R                  45      (       a  SnGOU R                    Vs/ s H  n[        XG5      (       d  M  UPM     nn[#        U5      S:X  d   SUR                   35       eUS   n	[        R
                  R%                  5       n
S/U
-  n[        R
                  R'                  5        Vs/ s H  n[(        U	   UR*                  :X  d  M  UPM      nn[#        U5      S:X  d   S5       e[-        U R/                  US   5      5      X½S   R0                  '   [3        U5      n[5        X&5      nGMí     Uc   eU$ s  snf s  snf )Nr‚   rN   úAmbiguous type: r   Ú1z# of Match expected to 1)r�   r   r   ÚTMPrL   ÚkernelÚcseÚvarname_mapÚnameÚshapeÚUNBACKED_INTÚSIZEÚPRECOMPUTED_SIZEÚINDEXÚFLOATÚUNBACKED_FLOATÚblock_typesr�   Útriton_tensor_ndimÚactive_range_treesr   ÚprefixÚstrÚget_block_sizeÚ
tensor_dimÚtupler:   )rŸ   r’   Ú
expr_shapeÚ	expr_varsÚvarÚcse_varÚ	var_shapeÚsymtÚsymbol_matchesÚsymÚndimrÉ   ÚtreeÚ
tree_matchs                 r†   Úget_block_shapeÚTritonSymbols.get_block_shapeÒ   sË  € ð &(ˆ
Ø×%Ñ%ˆ	ÜˆCÜ˜c¤4§8¡8×,Ñ,ÜŸ(™(Ÿ,™,×2Ñ2°3·8±8Ñ<�Ø#ŸM™M’	ÜØä×%Ñ%Ü—I‘IÜ×)Ñ)Ü—J‘JÜ—J‘JÜ×'Ñ'ð÷
ñ 
ð ’	ð &)§_¢_ó"Ú%4˜T¼Às×8Q—D¡_ð ð "ô ˜>Ó*¨aÓ/ÐNÐ3CÀCÇHÁHÀ:Ð1NÓNÐ/à$ QÑ'�Ü—x‘x×2Ñ2Ó4�Ø˜ ™�ô !"§¡× ;Ñ ;Ô =óâ =˜Ü! #‘¨$¯+©+Ñ5÷ Ù =ð ð ô
 ˜:“¨!Ó+ÐGÐ-GÓGÐ+ä25°c×6HÑ6HÈÐTUÉÓ6WÓ2X� ‘m×.Ñ.Ñ/Ü! %›L�	ô /¨zÓE‹JñK ðN Ñ%Ð%Ð%àÐùò1"ùòs   Ã0HÄHÅ<HÆHc                ó4   • U R                   UR                     $ rš   )Úblock_sizesrÝ   ©rŸ   rá   s     r†   rÕ   ÚTritonSymbols.get_block_size	  s   € à�‰˜tŸy™yÑ)Ð)rˆ   c                ó4   • U R                   UR                     $ rš   )Úblock_offsetsrÝ   rç   s     r†   Úget_block_offsetÚTritonSymbols.get_block_offset  s   € à× Ñ  §¡Ñ+Ð+rˆ   r‚   N)r’   ú
sympy.Exprr¦   rr   )rá   ra   r¦   zsympy.Symbol)r›   r¨   r©   rª   r«   r   r   ÚR0_INDEXÚR1_INDEXÚreduction_typesÚXBLOCKÚYBLOCKÚZBLOCKrÐ   r‹   ÚSymbolr   rê   Úupperræ   r­   rã   rÕ   rë   r®   )Ú.0rÝ   r‹   r   s   0000r†   r¸   r¸   ¾   s  † ññ ! $§-¡-°·±Ð!?Ó@€OÙ˜dŸk™k¨4¯;©;¸¿¹ÐVÀoÐVÓW€Kñ  öâˆDð 	Œe�lŠlœj¨Ñ.Ð/¨vÐ6ÀÐRVÑWÒWÙô€Mñ  ö	ò  ˆDð 	Œe�lŠlÜ˜$Ñ×%Ñ%Ó'Ð(¨Ð.¸Àtñ
ò 	
ñ  ô	€Kð ó4ó ð4ðl ó*ó ð*ð ó,ó ó,ùõQùõ
s   Á+C1
Â9C9
r¸   c                  ó˜   • \ rS rSr% S\S'   S\S'   S\S'   S\S	'   S
\S'   S\S'   SS jrSS jrSS jrSS jrSS jr	\
SS j5       rSrg)ÚIndexingOptionsi  rÔ   Ú	index_strúOrderedSet[str]Ú	mask_varszOptional[str]Ú
expand_strr¥   Ú_has_rindexrí   Úindexz#Optional[Sequence[Union[int, str]]]Úexpand_shapec                ó,   • [        U R                  5      $ rš   )r¥   rû   ©Úselfs    r†   Úhas_maskÚIndexingOptions.has_mask  s   € Ü�D—N‘NÓ#Ð#rˆ   c                óJ   • [        U R                  [        R                  5      $ rš   )r   rþ   r   rÄ   r  s    r†   Úhas_indirectÚIndexingOptions.has_indirect  s   € Ü" 4§:¡:¬t¯x©xÓ8Ð8rˆ   c                ó   • U R                   $ rš   )rý   r  s    r†   Ú
has_rindexÚIndexingOptions.has_rindex!  s   € Ø×ÑÐrˆ   c                ó:   • [        S U R                   5       5      $ )Nc              3  óV   #   • U  H  n[        U5      R                  S 5      v •  M!     g7f)ÚtmpN©rÔ   Ú
startswith©rö   Úmasks     r†   Ú	<genexpr>Ú.IndexingOptions.has_tmpmask.<locals>.<genexpr>%  s"   é € ÐJº>°4”3�t“9×'Ñ'¨×.Ð.º>ùó   ‚')©rz   rû   r  s    r†   Úhas_tmpmaskÚIndexingOptions.has_tmpmask$  s   € ÜÑJ¸4¿>º>ÓJÓJÐJrˆ   c                ó:   • [        S U R                   5       5      $ )Nc              3  óV   #   • U  H  n[        U5      R                  S 5      v •  M!     g7f)ÚrNr  r  s     r†   r  Ú,IndexingOptions.has_rmask.<locals>.<genexpr>(  s"   é € ÐHº°”3�t“9×'Ñ'¨×,Ð,ºùr  r  r  s    r†   Ú	has_rmaskÚIndexingOptions.has_rmask'  s   € ÜÑH¸¿ºÓHÓHÐHrˆ   c                óŒ   • U R                   (       a2  SR                  [        [        [        U R                   5      5      5      $ S$ )Nú & r§   )rû   ÚjoinÚsortedÚmaprÔ   r  s    r†   Úmask_strÚIndexingOptions.mask_str*  s4   € ð =A¿N¿NˆE�J‰J”vœc¤# t§~¡~Ó6Ó7Ó8ð	
ØPVð	
rˆ   r‚   N©r¦   r¥   ©r¦   rÔ   )r›   r¨   r©   rª   r¬   r  r  r	  r  r  Úpropertyr#  r®   r‚   rˆ   r†   rø   rø     sT   ‡ àƒNØÓØÓØÓØÓØ5Ó5ô$ô9ô ôKôIð ó
ó ó
rˆ   rø   c                  ó¨  • \ rS rSr% SrS\S'   S\S'   S\S'   S	\S
'   S\S'   S\S'   S\S'   S\S'   SrS\S'   SrS\S'   \S)S j5       r	\S)S j5       r
\S)S j5       r\S)S j5       r\SS.               S*S jj5       r        S+S jrS,S jr      S-S  jrS.S! jrS/S" jrS/S# jrS/S$ jrS/S% jrS/S& jr            S0S' jrS(rg)1ÚBlockDescriptorOptionsi4  zÎ
This is a base class that describes a block descriptor used in Triton kernels.
It can be used to create either a tensor descriptor (with TensorDescriptorOptions)
or a block pointer (with BlockPtrOptions).
ÚBlockParametersÚparamsrí   Úconstant_offsetú	list[int]Úorderrú   rû   úSequence[sympy.Expr]Úbroadcast_shapeú
list[bool]Úbroadcasting_dimsÚfinal_shapeúBlockParameters.StrideSorterÚstride_sorterNzOptional[list[int]]Ú_boundary_checkFr¥   Úcan_liftc                ó.   • U R                   R                  $ rš   )r+  rÉ   r  s    r†   rÉ   ÚBlockDescriptorOptions.shapeL  s   € à�{‰{× Ñ Ð rˆ   c                ó.   • U R                   R                  $ rš   )r+  Úblock_shaper  s    r†   r;  Ú"BlockDescriptorOptions.block_shapeP  s   € à�{‰{×&Ñ&Ð&rˆ   c                ó.   • U R                   R                  $ rš   )r+  Ústridesr  s    r†   r>  ÚBlockDescriptorOptions.stridesT  ó   € à�{‰{×"Ñ"Ð"rˆ   c                ó.   • U R                   R                  $ rš   )r+  Úoffsetsr  s    r†   rB  ÚBlockDescriptorOptions.offsetsX  r@  rˆ   )r7  c               ó,  ^• [         R                  R                  mS	U4S jjnU" UR                  5      Ul        U" UR                  5      Ul        UR
                   V	s/ s H  n	TR                  U	S5      PM     n
n	[        U
5      (       a  SU
S'   UR                  U
5      nUR                  U[         R                  R                  S9u  pUR                   Vs/ s H  nTR                  US5      PM     nnUR
                  nUR                  U5      nU Vs/ s H  n[        R                  U5      PM     nn[         R                  R                  (       a&  US   R                  S:X  d   eUR!                  S5        [         R                  R"                  n[         R                  R$                  (       d‹  ['        UR                  5      ['        [         R                  R(                  5      U-
  :X  aN  [         R                  R*                  R-                  5       (       a!  U[.        R0                  R2                  /U-  -  n [4        R6                  " [         R                  R                  UR                  5      nU " U[         R                  R                  RA                  U5      UUUUUUUS9	nURC                  XS5        U$ s  sn	f s  snf s  snf ! [8         a3    [;        [=        [?        ['        UR                  5      5      5      5      n N“f = f)
z2Helper to create a BlockDescriptorOptions instancec                óR   >• U  Vs/ s H  nTR                  U5      PM     sn$ s  snf rš   )Úlookup_precomputed_size)Úexprsr’   Úsizevarss     €r†   Úlookup_sizeÚ2BlockDescriptorOptions.create.<locals>.lookup_sizel  s&   ø€ ÙGLÓMÂu¸t�H×4Ñ4°TÖ:ÁuÑMÐMùÒMs   †$rN   Féÿÿÿÿ)Ústride_sorter_clsÚ	shape_envr   Úx)	r+  r,  r.  rû   r3  r0  r2  r5  r7  )rG  zIterable[sympy.Expr]r¦   úlist[sympy.Expr])"rL   ÚgraphrH  rÉ   r>  r;  Ústatically_known_equalsÚallÚremove_dimsÚmaybe_sort_with_stride_orderÚ
_shape_envr¸   rÕ   rÅ   Úno_x_dimrÓ   ÚpopÚnum_reduction_dimsÚinside_reductionr�   ÚnumelsÚfeaturesÚis_reductionr‹   ÚSÚOner%   Úargsort_symÚAssertionErrorÚlistÚreversedÚrangerF  Úcompute_boundary_check)rŸ   r+  r,  Úrange_treesrû   Úget_max_blockrL  r7  rI  ÚdimÚsingleton_dimsr5  Ústrider2  r0  rá   r3  Úreduction_ndimr.  ÚresultrH  s                       @r†   ÚcreateÚBlockDescriptorOptions.create\  s’  ø€ ô —7‘7×#Ñ#ˆ÷	Nñ # 6§<¡<Ó0ˆŒÙ$ V§^¡^Ó4ˆŒð AG×@RÒ@Ró
Ú@R¸ˆH×,Ñ,¨S°!Ö4Ñ@Rð 	ð 
ô ˆ~×Ñà!&ˆN˜2Ñð ×#Ñ# NÓ3ˆð !'× CÑ CØ/¼1¿7¹7×;MÑ;Mð !Dð !
Ñˆð GMÇnÂnó
ÚFT¸FˆH×,Ñ,¨V°QÖ7Ánð 	ð 
ð !×,Ñ,ˆð ×#Ñ#Ð$5Ó6ˆñ GRÓRÂk¸d”}×3Ñ3°DÖ9ÁkˆÐRÜ�8‰8××Ø˜q‘>×(Ñ(¨CÓ/Ð/Ð/Ø�O‰O˜AÔäŸ™×4Ñ4ˆä—‘×)×)Ü�F—N‘NÓ#¤s¬1¯8©8¯?©?Ó';¸nÑ'LÓLÜ—‘×!Ñ!×.Ñ.×0Ñ0ð œEŸG™GŸK™K˜=¨>Ñ9Ñ9ˆKð	?ô ×%Ò%¤a§g¡g×&8Ñ&8¸&¿.¹.ÓIˆEñ
 ØÜŸG™G×,Ñ,×DÑDÀ_ÓUØØØ#Ø+Ø/Ø'Øñ

ˆð 	×%Ñ% mÔAØˆùòC
ùò$
ùò Søô$ ó 	?äœ¤%¬¨F¯N©NÓ(;Ó"<Ó=Ó>ŠEð	?ús$   Á KÃKÄKÉ9K Ë:LÌLc                óB   • [         R                  U   n[        XU05      $ ©z>
Replaces instances of {symt}_offset with the new expression.
©r¸   rê   rE   ©r  r’   ÚreplacementrÝ   Úroffsets        r†   Úreplace_offsetÚ%BlockDescriptorOptions.replace_offsetº  ó$   € ô  ×-Ñ-¨dÑ3ˆÜ˜$¨+Ð 6Ó7Ð7rˆ   c                ó‚   • [         R                   H*  nU R                  U[        R                  " S5      U5      nM,     U$ ©Nr   ©r¸   rð   rt  r‹   r�   ©r  r’   rÝ   s      r†   Úremove_roffsetsÚ&BlockDescriptorOptions.remove_roffsetsÃ  ó4   € Ü!×1Ô1ˆDØ×&Ñ& t¬U¯]ª]¸1Ó-=¸tÓDŠDñ 2àˆrˆ   c           
     ó&  • [         R                  R                  nU Vs0 s H8  n[        R                  UR
                     U" [        UR
                     5      _M:     nn[        [        [         R                  R                  U5      5      n[        [        U R                  5      5       Vs/ s GHK  nUR                  U R                  U   [         R"                  R$                  5      (       a  MB  U(       a<  [        R                  [&        R(                     U R*                  U   R,                  ;   dp  UR/                  U R                  U   U R*                  U   5      (       a  M¸  UR/                  U R                  U   [1        U R*                  U   U5      5      (       a  Mõ  [         R                  R2                  (       a5  U R*                  U   [        R                  [&        R4                     :X  a  GMI  UPGMN     snU l        gs  snf s  snf )z6List of indices to pass to tl.load(boundary_check=...)N)rL   rP  rH  r¸   ræ   rÝ   r   rz   r"  rÅ   Úneeds_yz_grid_overflowrc  r�   rÉ   rQ  r>  r‹   r]  ÚZeror   rò   r;  r�   Ústatically_known_multiple_ofrE   rV  rñ   r6  )r  rf  re  rH  ÚtÚblock_to_maxÚneeds_overflow_gridÚidxs           r†   rd  Ú-BlockDescriptorOptions.compute_boundary_checkÈ  s™  € ô —7‘7×#Ñ#ˆñ !ó/
â �ô ×%Ñ% a§f¡fÑ-©}¼ZÈÏÉÑ=OÓ/PÒPÙ ð 	ð /
ô "¤#¤a§h¡h×&EÑ&EÀ{Ó"SÓTÐô œS §¡›_Ô-ó 
ã-�à×4Ñ4°T·\±\À#Ñ5FÌÏÉÏÉ×Uó ö ,Ü)×5Ñ5´d·k±kÑBØ×+Ñ+¨CÑ0×=Ñ=ó>ð %×AÑAØ ŸJ™J s™O¨T×-=Ñ-=¸cÑ-B÷ó ð !)× EÑ EØ ŸJ™J s™OÜ& t×'7Ñ'7¸Ñ'<¸lÓK÷!ó ô* —H‘H×%×%Ø×(Ñ(¨Ñ-´×1JÑ1JÌ4Ï;É;Ñ1WÑW÷- ‰CÙ-ñ 
ˆÕùò/
ùò 
s&   Ÿ?H	Â-AHÃ2A2HÅ(9HÆ%AHÇ9Hc                ó8   • U R                   c   eU R                   $ rš   ©r6  r  s    r†   Úboundary_checkÚ%BlockDescriptorOptions.boundary_checkü  s    € Ø×#Ñ#Ñ/Ð/Ð/Ø×#Ñ#Ð#rˆ   c                ó   • g©NFr‚   r  s    r†   r  Ú#BlockDescriptorOptions.has_indirect   ó   € Ørˆ   c                ó:   • [        S U R                   5       5      $ )Nc              3  óV   #   • U  H  n[        U[        R                  5      v •  M!     g 7frš   )r   r¸   rð   )rö   r’   s     r†   r  Ú4BlockDescriptorOptions.has_rindex.<locals>.<genexpr>  s'   é € ð 
â(�ô   ¤m×&CÑ&C×DÐDÚ(ùr  )rz   r;  r  s    r†   r	  Ú!BlockDescriptorOptions.has_rindex  s"   € Üñ 
à×(Ò(ó
ó 
ð 	
rˆ   c                ó"   • U R                  5       $ rš   )r	  r  s    r†   r  Ú BlockDescriptorOptions.has_rmask	  s   € Ø�‰Ó Ð rˆ   c                ó   • grŒ  r‚   r  s    r†   r  Ú"BlockDescriptorOptions.has_tmpmask  rŽ  rˆ   c                ó4   • [        U R                  5       5      $ rš   )r¥   r‰  r  s    r†   r  ÚBlockDescriptorOptions.has_mask  s   € Ü�D×'Ñ'Ó)Ó*Ð*rˆ   c                óÂ  ^• U R                   nU R                  nU(       aJ  U R                  R                  U R                   5      nU R                  R                  U R                  5      n[	        Xg5       VV	s/ s H(  u  p‰U	(       a  [
        R                  R                  OUPM*     n
nn	[        XU
5      nU R                  R                  (       d:  U(       d3  [        U
5      [        U5      :X  a  U R                  R                  U
5      n
[        R                  R                  mU=(       a<    [        U
5      [        U5      :H  =(       a    [        U4S j[	        X£5       5       5      n[        U R                  5      (       a-  U(       d&  SU S[        R                   R#                  U5       S3nU R                   nU R                  R                  (       dt  U(       a  U R                  R$                  OU R                  R&                  nSU SU S3nU(       a  U R                   O$U R                  R                  U R                   5      n[        XU5      nU$ s  sn	nf )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              3  ó|   >#   • U  H1  u  pTR                  US 5      =(       d    TR                  X5      v •  M3     g7f©rN   N)rQ  )rö   Úpre_dimÚpost_dimrH  s      €r†   r  ÚGBlockDescriptorOptions.codegen_broadcast_and_reshape.<locals>.<genexpr>P  sE   øé € ð ò *OÑ%�Gð ×0Ñ0°¸!Ó<÷ GØ×3Ñ3°GÓFôGâ)Nùs   ƒ9<útl.broadcast_to(ú, Ú)ú	tl.trans()r0  r2  r5  ÚrevertÚzipr‹   r]  r^  Útriton_reshapeÚis_identityr�   rL   rP  rH  rR  rz   rÅ   Úindex_to_strÚsort_idxÚrevert_sort_idx)r  ÚvalueÚinitial_shaper3  Úallow_implicitÚ	for_storer0  r2  rg  Úis_broadcastingÚpre_broadcast_shapeÚsupports_implicit_broadcastÚ	old_shapeÚpermute_dimsrH  s                 @r†   Úcodegen_broadcast_and_reshapeÚ4BlockDescriptorOptions.codegen_broadcast_and_reshape  s  ø€ ð2 ×.Ñ.ˆØ ×2Ñ2Ðö Ø"×0Ñ0×7Ñ7¸×8LÑ8LÓMˆOØ $× 2Ñ 2× 9Ñ 9¸$×:PÑ:PÓ QÐô
 ),¨OÔ(Oô
â(OÑ$�ö +ŒE�G‰G�KŠK°Ò3Ù(Oð 	ñ 
ô ˜uÐ5HÓIˆð ×"Ñ"×.×.ÞÜÐ'Ó(¬C°Ó,<Ó<ð #'×"4Ñ"4×";Ñ";Ð<OÓ"PÐô —7‘7×#Ñ#ˆØ&4÷ '
ÜÐ#Ó$¬¨KÓ(8Ñ8÷ Üô ô *-Ð-@Ô)Nóó ð 	$ô ˆt×%Ñ%×&Ñ&Ö/Jà" 5 '¨¬A¯H©H×,AÑ,AÀ/Ó,RÐ+SÐSTÐUð ð ×(Ñ(ˆ	Ø×!Ñ!×-×-ö ð ×"Ñ"×+Ò+à×'Ñ'×7Ñ7ð ð
   ˜w b¨¨°aÐ8ˆEö ð ×$Ò$à×'Ñ'×.Ñ.¨t×/CÑ/CÓDð ô ˜u°Ó=ˆàˆùów
s   Á9/Irˆ  ©r¦   rO  )r+  r*  r,  rí   re  úlist[IterationRangesRoot]rû   rú   rf  úCallable[[str], int]rL  z"type[BlockParameters.StrideSorter]r7  r¥   r¦   r)  ©r’   rí   rr  rí   rÝ   r   r¦   rí   ©r’   rí   r¦   rí   )rf  r·  re  r¶  r¦   r§   )r¦   r-  r%  )rª  rÔ   r«  r/  r3  r/  r¬  r¥   r­  r¥   r¦   rÔ   )r›   r¨   r©   rª   r«   r¬   r6  r7  r'  rÉ   r;  r>  rB  r­   rl  rt  r{  rd  r‰  r  r	  r  r  r  r³  r®   r‚   rˆ   r†   r)  r)  4  sÄ  ‡ ñð ÓØÓØÓØÓØ)Ó)Ø!Ó!Ø%Ó%ð 0Ó/Ø+/€OÐ(Ó/ð €HˆdÓàó!ó ð!ð ó'ó ð'ð ó#ó ð#ð ó#ó ð#ð ð ñ[ð  ð[ð $ð	[ð
 /ð[ð #ð[ð ,ð[ð >ð[ð ð[ð 
 ô[ó ð[ðz8Øð8Ø-7ð8Ø?Cð8à	ô8ôð
2
à+ð2
ð /ð2
ð 
ô	2
ôh$ôô
ô!ôô+ðaàðað ,ðað *ð	að
 ðað ðað 
÷arˆ   r)  c                  ó"   • \ rS rSrSSS jjrSrg)ÚTensorDescriptorOptionsiv  c                ó.  • [         R                  R                  nU R                  S:w  a  U SU" U R                  5       S3OUSU" U R                  5       3SU" U R
                  5       3SU" U R                  5       3/nSSR                  U5       S3$ )	z×
Codegen a call to tl.make_tensor_descriptor()

Args:
    name: variable name for pointer
    roffset: unused, but kept for compatibility with BlockPtrOptions.format()

Returns:
    "tl.make_tensor_descriptor(...)"
r   ú + (r¡  úshape=ústrides=úblock_shape=ztl.make_tensor_descriptor(r   )rL   rÅ   r§  r,  rÉ   r>  r;  r   )r  rÈ   rs  ÚfÚargss        r†   ÚformatÚTensorDescriptorOptions.formatx  s    € ô �H‰H×!Ñ!ˆð ×'Ñ'¨1Ó,ð �&˜™Q˜t×3Ñ3Ó4Ð5°QÑ7àà‘Q�t—z‘z“]�OÐ$Ø‘q˜Ÿ™“Ð'Ð(Ø™1˜T×-Ñ-Ó.Ð/Ð0ð	
ˆð ,¨D¯I©I°d«OÐ+<¸AÐ>Ð>rˆ   r‚   N©T©rÈ   rÔ   r¦   rÔ   )r›   r¨   r©   rª   rÃ  r®   r‚   rˆ   r†   r»  r»  v  s   † ÷?ñ ?rˆ   r»  c                  óP   • \ rS rSr        SS jrS	S jrS
SS jjrSS jrSrg)ÚBlockPtrOptionsi“  c                óB   • [         R                  U   n[        XU05      $ ro  rp  rq  s        r†   rt  ÚBlockPtrOptions.replace_offset•  rv  rˆ   c                ó‚   • [         R                   H*  nU R                  U[        R                  " S5      U5      nM,     U$ rx  ry  rz  s      r†   r{  ÚBlockPtrOptions.remove_roffsetsž  r}  rˆ   c           	     óÞ  • [         R                  R                  n/ U R                  QnU(       d   U Vs/ s H  oPR	                  U5      PM     nnU R
                  S:w  a  U SU" U R
                  5       S3OUSU" U R                  5       3SU" U R                  5       3SU" U R                  5       3SU" U R                  5       3SU" U5       3/nS	S
R                  U5       S3$ s  snf )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(...)"
r   r½  r¡  r¾  r¿  rÀ  zorder=zoffsets=ztl.make_block_ptr(r   )rL   rÅ   r§  rB  r{  r,  rÉ   r>  r;  r.  r   )r  rÈ   rs  rÁ  rB  rº   rÂ  s          r†   rÃ  ÚBlockPtrOptions.format£  só   € ô �H‰H×!Ñ!ˆØ!�D—L‘L�/ˆÞÙBIÓJÂ'¸×+Ñ+¨FÖ3Á'ˆGÐJð ×'Ñ'¨1Ó,ð �&˜™Q˜t×3Ñ3Ó4Ð5°QÑ7àà‘Q�t—z‘z“]�OÐ$Ø‘q˜Ÿ™“Ð'Ð(Ø™1˜T×-Ñ-Ó.Ð/Ð0Ø‘Q�t—z‘z“]�OÐ$Ø‘q˜“z�lÐ#ð
ˆð $ D§I¡I¨d£OÐ#4°AÐ6Ð6ùò Ks   ´C*c           	     óæ   • [         R                  U   nU R                   Vs/ s HA  nU R                  X2U5      U R                  U[        R
                  R                  U5      -
  PMC     nnU$ s  snf )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.
)r¸   ræ   rB  rt  r‹   r]  r€  )r  rÝ   Úrblockrº   Úadvances        r†   Úadvance_roffsetÚBlockPtrOptions.advance_roffsetÀ  sw   € ô ×*Ñ*¨4Ñ0ˆð Ÿ,š,ó
ò
 '�ð ×#Ñ# F°DÓ9Ø×%Ñ% f¬e¯g©g¯l©l¸DÓAôBñ 'ð 	ð 
ð ˆùò
s   ¢AA.r‚   Nr¸  r¹  rÅ  rÆ  )rÝ   r   r¦   rí   )	r›   r¨   r©   rª   rt  r{  rÃ  rÒ  r®   r‚   rˆ   r†   rÈ  rÈ  “  s6   † ð8Øð8Ø-7ð8Ø?Cð8à	ô8ôö
7÷:rˆ   rÈ  c                ó¨  • [        U[        5      (       a  [        U[        5      (       d   eU Vs/ s H"  n[        R                  R	                  U5      PM$     nnU Vs/ s H"  n[        R                  R	                  U5      PM$     nnXE:X  a  U $ U Vs/ s H  ofS:w  d  M
  UPM     snU:w  a  SU  SSR                  U5       S3$ Sn/ nU HK  n	U[        U5      :  a   X”U   :X  a  UR                  S5        US-  nM2  U	S:X  d   eUR                  S	5        MM     U[        U5      :X  d   eU  S
SR                  U5       S3$ s  snf s  snf s  snf )z<Workaround https://github.com/triton-lang/triton/issues/2836rÃ   ztl.reshape(z, [r   z])r   Ú:rN   r§   Ú[Ú])rŠ   ra  rL   rÅ   r§  r   r�   Úappend)
rª  r±  Ú	new_shaperÉ   Úold_shape_strÚnew_shape_strÚsr…  ÚexpandÚsizes
             r†   r¥  r¥  Ô  sF  € ô �i¤×&Ñ&¬:°iÄ×+FÑ+FÐFÐFá?HÓIºy°e”Q—X‘X×*Ñ*¨5Ö1¹y€MÐIÙ?HÓIºy°e”Q—X‘X×*Ñ*¨5Ö1¹y€MÐIàÓ%ØˆÙ Ó-’=ˆa¨¡H�‘=Ñ-°Ó>Ø˜U˜G 3 t§y¡y°Ó'?Ð&@ÀÐCÐCà
€CØ€FÛˆØ”�]Ó#Ó#¨¸cÑ0BÓ(BØ�M‰M˜#ÔØ�1‰HŠCà˜3“;Ð�;Ø�M‰M˜&Ö!ñ ð ”#�mÓ$Ó$Ð$Ð$ØˆW�A�d—i‘i Ó'Ð(¨Ð*Ð*ùò% JùÚIùò .s   ±)EÁ )E
Â	EÂ#Ec                  ó@  • \ rS rSrS"S jrS"S jrS"S jrS"S jrS"S jrS"S jr	S"S jr
S"S	 jrS"S
 jrS"S jrS"S jrS"S jrS"S jrS"S jrS#S jrS"S jrS"S jrS"S jrS"S jrS"S jrS"S jrS"S jrS"S jrS"S jrS"S jrS"S jrS"S jrS"S jrS"S jr S"S jr!S r"g!)$ÚTritonPrinterió  c                ó²   • [        UR                  5      S:X  d   eSU R                  UR                  S   5       S[        R                  R
                   S3$ )NrN   úlibdevice.trunc(r   ú).to(r¡  ©r�   rÂ  Ú_printrL   rÅ   Úindex_dtype©r  r’   s     r†   Ú_print_TruncToIntÚTritonPrinter._print_TruncToIntô  óO   € Ü�4—9‘9‹~ Ó"Ð"Ð"ð ˜tŸ{™{¨4¯9©9°Q©<Ó8Ð9¸¼q¿x¹x×?SÑ?SÐ>TÐTUÐVð	
rˆ   c                óÜ   • UR                   (       a  [        [        U5      5      nU$ [        R                  " 5       (       a$  [
        R                  R                  (       a  U nU$ SU S3nU$ )Nútl.full([], ú, tl.float64))rŽ   rÔ   Úintr"   Ú	is_fbcoderœ   ÚversionÚhip)r  r’   Úrets      r†   Ú_print_FloatÚTritonPrinter._print_Floatû  s_   € Ø�?�?ô ”c˜$“i“.ˆCð
 ˆ
ô	 ×Ò×Ñ¤E§M¡M×$5×$5Ø�FˆCð ˆ
ð !   mÐ4ˆCØˆ
rˆ   c                ó”   • [        UR                  5      S:X  d   eU R                  UR                  S   [        S   S-
  5      nU S3$ )NrN   r   ÚAtomç      à?z.to(tl.float64))r�   rÂ  Úparenthesizer   )r  r’   rÜ  s      r†   Ú_print_ToFloatÚTritonPrinter._print_ToFloat  sI   € Ü�4—9‘9‹~ Ó"Ð"Ð"à×Ñ˜dŸi™i¨™l¬J°vÑ,>ÀÑ,DÓEˆØ��OÐ$Ð$rˆ   c                ó  • UR                   u  p#UR                  (       a8  UR                  (       a'  U R                  UR                   S[        S   S-
  5      $ U R	                  U5      nU R	                  U5      nSU SU S3$ )Nú % rö  r÷  z!triton_helpers.remainder_integer(r   r¡  )rÂ  Úis_nonnegativeÚ	stringifyr   rå  ©r  r’   ÚquotÚdivÚquot_sÚdiv_ss         r†   Ú_print_PythonModÚTritonPrinter._print_PythonMod  sp   € Ø—I‘I‰	ˆØ×× 3×#5×#5Ø—>‘> $§)¡)¨U´J¸vÑ4FÈÑ4LÓMÐMà—‘˜TÓ"ˆà—‘˜CÓ ˆØ2°6°(¸"¸U¸GÀ1ÐEÐErˆ   c                ó,  • UR                   (       d   eUR                  u  p#UR                  (       a8  UR                  (       a'  U R                  UR                  S[        S   S-
  5      $ U R                  U5      nU R                  U5      nSU SU S3$ )Nú // rö  r÷  z!triton_helpers.div_floor_integer(z,  r¡  )rŽ   rÂ  rý  rþ  r   rå  rÿ  s         r†   Ú_print_FloorDivÚTritonPrinter._print_FloorDiv  s|   € Ø��ÐˆØ—I‘I‰	ˆØ×× 3×#5×#5Ø—>‘> $§)¡)¨V´ZÀÑ5GÈ#Ñ5MÓNÐNà—‘˜TÓ"ˆà—‘˜CÓ ˆØ2°6°(¸#¸e¸WÀAÐFÐFrˆ   c                óP   • U R                  UR                  S[        S   S-
  5      $ )Nú / rö  r÷  )rþ  rÂ  r   rç  s     r†   Ú_print_IntTrueDivÚTritonPrinter._print_IntTrueDiv$  s#   € Ø�~‰~˜dŸi™i¨´
¸6Ñ0BÀSÑ0HÓIÐIrˆ   c                ó²   • [        UR                  5      S:X  d   eSU R                  UR                  S   5       S[        R                  R
                   S3$ ©NrN   úlibdevice.floor(r   rã  r¡  rä  rç  s     r†   Ú_print_floorÚTritonPrinter._print_floor)  rê  rˆ   c                ó²   • [        UR                  5      S:X  d   eSU R                  UR                  S   5       S[        R                  R
                   S3$ r  rä  rç  s     r†   Ú_print_FloorToIntÚTritonPrinter._print_FloorToInt0  rê  rˆ   c                ó²   • [        UR                  5      S:X  d   eSU R                  UR                  S   5       S[        R                  R
                   S3$ ©NrN   úlibdevice.ceil(r   rã  r¡  rä  rç  s     r†   Ú_print_ceilingÚTritonPrinter._print_ceiling7  óK   € Ü�4—9‘9‹~ Ó"Ð"Ð"à  §¡¨T¯Y©Y°q©\Ó!:Ð ;¸5ÄÇÁ×AUÑAUÐ@VÐVWÐXÐXrˆ   c                ó²   • [        UR                  5      S:X  d   eSU R                  UR                  S   5       S[        R                  R
                   S3$ r  rä  rç  s     r†   Ú_print_CeilToIntÚTritonPrinter._print_CeilToInt<  r  rˆ   c                ó,   • SU R                  U5       S3$ )Nztl.sqrt_rn((ú).to(tl.float32)))rå  rç  s     r†   Ú_helper_sqrtÚTritonPrinter._helper_sqrtA  s   € à˜dŸk™k¨$Ó/Ð0Ð0AÐBÐBrˆ   c                óŒ   • U R                  UR                  S   5      nU R                  UR                  S   5      nSU SU S3$ )Nr   rN   zlibdevice.pow((z).to(tl.float64), (z).to(tl.float64)))rå  rÂ  )r  r’   ÚbaseÚexps       r†   Ú_print_FloatPowÚTritonPrinter._print_FloatPowE  sG   € à�{‰{˜4Ÿ9™9 Q™<Ó(ˆà�k‰k˜$Ÿ)™) A™,Ó'ˆð !  Ð&9¸#¸Ð>OÐPÐPrˆ   c                óX  • UR                   S   R                  (       a  S[        UR                   S   5       S3nO"SU R                  UR                   S   5       S3nUR                   S   nUR                  (       a  S[        U5       S3nOSU R                  U5       S3nSU SU S	3$ )
Nr   rì  rí  Ú(z).to(tl.float64)rN   úlibdevice.pow(r   r¡  )rÂ  Ú
is_IntegerÚfloatrå  )r  r’   r$  Úexp_valr%  s        r†   Ú_print_PowByNaturalÚ!TritonPrinter._print_PowByNaturalO  s­   € Ø�9‰9�Q‰<×"×"Ø!¤%¨¯	©	°!©Ó"5Ð!6°mÐD‰Dð �t—{‘{ 4§9¡9¨Q¡<Ó0Ð1Ð1AÐBˆDØ—)‘)˜A‘,ˆØ××Ø ¤ w£Ð 0°Ð>‰Cð �d—k‘k 'Ó*Ð+Ð+;Ð<ˆCð   ˜v R¨ u¨AÐ.Ð.rˆ   c                óÎ   • U R                  UR                  S   5      nU R                  UR                  S   5      nU R                  UR                  S   5      nSU SU SU S3$ )Nr   rN   r!   ú	tl.where(r   r¡  )ÚdoprintrÂ  )r  r’   ÚcÚpÚqs        r†   Ú_print_WhereÚTritonPrinter._print_Where`  s_   € Ø�L‰L˜Ÿ™ 1™Ó&ˆØ�L‰L˜Ÿ™ 1™Ó&ˆØ�L‰L˜Ÿ™ 1™Ó&ˆØ˜1˜#˜R ˜s " Q C qÐ)Ð)rˆ   c                ó¼  • [        UR                  5      S:X  a  U R                  UR                  S   5      $ [        UR                  5      S-  n[        U5      nU R                  U" UR                  SU 6 5      nU R                  U" UR                  US 6 5      n[	        S XV4 5       5      u  pVUS;   d   SU S35       eS	U S
U SU SU SU S
U SU SU S3$ )z1
Helper for max/min code generation.
cmp: > or <
rN   r   r!   Nc              3  ó.   #   • U  H  nS U S3v •  M     g7f)r)  r¡  Nr‚   ©rö   rN  s     r†   r  Ú6TritonPrinter._print_min_max_helper.<locals>.<genexpr>x  s   é € Ð.¢v !�q˜˜˜1•X¢vùs   ‚)Ú>Ú<zUnexpected comparator: 'Ú'r)  z * (Ú z= z) + ú)))r�   rÂ  rå  Útyper×   )r  r’   ÚcmpÚmidrŸ   ÚaÚbs          r†   Ú_print_min_max_helperÚ#TritonPrinter._print_min_max_helperf  sì   € ô
 ˆt�y‰y‹>˜QÓà—;‘;˜tŸy™y¨™|Ó,Ð,ä�$—)‘)‹n Ñ!ˆÜ�4‹jˆà�K‰K™˜TŸY™Y t¨˜_Ð-Ó.ˆà�K‰K™˜TŸY™Y s t˜_Ð-Ó.ˆô Ñ.¨¡vÓ.Ó.‰ˆØ�jÓ ÐCÐ$<¸S¸EÀÐ"CÓCÐ Ø�1�#�T˜!˜˜A˜c˜U " Q C t¨A¨3¨d°1°#°Q°s°e¸1¸Q¸C¸rÐBÐBrˆ   c                ó&   • U R                  US5      $ )Nr=  ©rF  rç  s     r†   Ú
_print_MinÚTritonPrinter._print_Min|  ó   € Ø×)Ñ)¨$°Ó4Ð4rˆ   c                ó&   • U R                  US5      $ )Nr<  rI  rç  s     r†   Ú
_print_MaxÚTritonPrinter._print_Max  rL  rˆ   c                ó|   • [        UR                  5      S:X  d   eSU R                  UR                  S   5       S3$ )NrN   útl_math.abs(r   r¡  ©r�   rÂ  rå  rç  s     r†   Ú
_print_AbsÚTritonPrinter._print_Abs‚  s9   € Ü�4—9‘9‹~ Ó"Ð"Ð"à˜dŸk™k¨$¯)©)°A©,Ó7Ð8¸Ð:Ð:rˆ   c                ó|   • [        UR                  5      S:X  d   eSU R                  UR                  S   5       S3$ )NrN   zlibdevice.cos((r   r   rR  rç  s     r†   Ú_print_OpaqueUnaryFn_cosÚ&TritonPrinter._print_OpaqueUnaryFn_cos‡  ó:   € Ü�4—9‘9‹~ Ó"Ð"Ð"à  §¡¨T¯Y©Y°q©\Ó!:Ð ;Ð;LÐMÐMrˆ   c                ó|   • [        UR                  5      S:X  d   eSU R                  UR                  S   5       S3$ )NrN   zlibdevice.cosh((r   r   rR  rç  s     r†   Ú_print_OpaqueUnaryFn_coshÚ'TritonPrinter._print_OpaqueUnaryFn_coshŒ  ó:   € Ü�4—9‘9‹~ Ó"Ð"Ð"à! $§+¡+¨d¯i©i¸©lÓ";Ð!<Ð<MÐNÐNrˆ   c                ó|   • [        UR                  5      S:X  d   eSU R                  UR                  S   5       S3$ )NrN   zlibdevice.acos((r   r   rR  rç  s     r†   Ú_print_OpaqueUnaryFn_acosÚ'TritonPrinter._print_OpaqueUnaryFn_acos‘  r\  rˆ   c                ó|   • [        UR                  5      S:X  d   eSU R                  UR                  S   5       S3$ )NrN   zlibdevice.sin((r   r   rR  rç  s     r†   Ú_print_OpaqueUnaryFn_sinÚ&TritonPrinter._print_OpaqueUnaryFn_sin–  rX  rˆ   c                ó|   • [        UR                  5      S:X  d   eSU R                  UR                  S   5       S3$ )NrN   zlibdevice.sinh((r   r   rR  rç  s     r†   Ú_print_OpaqueUnaryFn_sinhÚ'TritonPrinter._print_OpaqueUnaryFn_sinh›  r\  rˆ   c                ó|   • [        UR                  5      S:X  d   eSU R                  UR                  S   5       S3$ )NrN   zlibdevice.asin((r   r   rR  rç  s     r†   Ú_print_OpaqueUnaryFn_asinÚ'TritonPrinter._print_OpaqueUnaryFn_asin   r\  rˆ   c                ó|   • [        UR                  5      S:X  d   eSU R                  UR                  S   5       S3$ )NrN   zlibdevice.tan((r   r   rR  rç  s     r†   Ú_print_OpaqueUnaryFn_tanÚ&TritonPrinter._print_OpaqueUnaryFn_tan¥  rX  rˆ   c                ó|   • [        UR                  5      S:X  d   eSU R                  UR                  S   5       S3$ )NrN   zlibdevice.tanh((r   r   rR  rç  s     r†   Ú_print_OpaqueUnaryFn_tanhÚ'TritonPrinter._print_OpaqueUnaryFn_tanhª  r\  rˆ   c                ó|   • [        UR                  5      S:X  d   eSU R                  UR                  S   5       S3$ )NrN   zlibdevice.atan((r   r   rR  rç  s     r†   Ú_print_OpaqueUnaryFn_atanÚ'TritonPrinter._print_OpaqueUnaryFn_atan¯  r\  rˆ   c                ó|   • [        UR                  5      S:X  d   eSU R                  UR                  S   5       S3$ )NrN   zlibdevice.log2((r   r   rR  rç  s     r†   Ú_print_OpaqueUnaryFn_log2Ú'TritonPrinter._print_OpaqueUnaryFn_log2´  r\  rˆ   c                ó²   • [        UR                  5      S:X  d   eSU R                  UR                  S   5       S[        R                  R
                   S3$ )NrN   zlibdevice.llrint(r   rã  r¡  rä  rç  s     r†   Ú_print_RoundToIntÚTritonPrinter._print_RoundToInt¹  sO   € Ü�4—9‘9‹~ Ó"Ð"Ð"ð   §¡¨D¯I©I°a©LÓ 9Ð:¸%ÄÇÁ×@TÑ@TÐ?UÐUVÐWð	
rˆ   c                óî   • [        UR                  5      S:X  d   eUR                  u  p#UR                  (       a  US:  d   e[        SU S35      eU R	                  U[
        S   5      nSU SU SU*  3$ )	Nr!   r   zOFor integer inputs, only non-negative ndigits are currently supported, but got r€   ÚMulzlibdevice.nearbyint(1eú * z) * 1e)r�   rÂ  rŽ   Ú
ValueErrorrø  r   )r  r’   ÚnumberÚndigitsÚ
number_strs        r†   Ú_print_RoundDecimalÚ!TritonPrinter._print_RoundDecimalÀ  sˆ   € Ü�4—9‘9‹~ Ó"Ð"Ð"ØŸ)™)‰ˆØ××à˜Q“;Ð�;ÜØaÐbiÐajÐjkÐlóð ð ×&Ñ& v¬z¸%Ñ/@ÓAˆ
Ø'¨ y°°J°<¸vÀwÀhÀZÐPÐPrˆ   r‚   N)r’   rí   r¦   rÔ   )r’   rí   rB  rÔ   r¦   rÔ   )#r›   r¨   r©   rª   rè  ró  rù  r  r  r  r  r  r  r  r!  r&  r.  r6  rF  rJ  rN  rS  rV  rZ  r^  ra  rd  rg  rj  rm  rp  rs  rv  r  r®   r‚   rˆ   r†   rà  rà  ó  s«   † ô
ô
ô%ôFô	GôJô

ô
ôYô
Yô
CôQô/ô"*ôCô,5ô5ô;ô
Nô
Oô
Oô
Nô
Oô
Oô
Nô
Oô
Oô
Oô

÷Qrˆ   rà  c                ó*   • [        [        U 5      5      $ )zCConvert torch.dtype to triton type and upcast [b]float16 to float32)rF   rH   ©Údtypes    r†   Útriton_compute_typer„  Ñ  s   € äÔ*¨5Ó1Ó2Ð2rˆ   c                ó`   • U [         R                  :X  a  [         R                  n [        U 5      $ )z@Convert torch.dtype to triton type, with fix for storing tl.bool)rœ   r¥   Úint8rF   r‚  s    r†   Útriton_store_typer‡  Ö  s"   € à”—
‘
ÓÜ—
‘
ˆÜ�uÓÐrˆ   c                óš   • [        U 5      (       a1  U R                  (       a   U R                  S::  a  [        R                  $ [        U 5      $ )z0Implicit upcasts used for Triton reduction typesé   )r   Ú	is_signedÚitemsizerœ   Úint32rH   r‚  s    r†   Úupcast_acc_dtyper�  Ý  s3   € ä˜×Ñ 5§?§?°u·~±~ÈÓ7JÜ�{‰{ÐÜ˜uÓ%Ð%rˆ   c                ó*   • [        [        U 5      5      $ )z:Convert torch.dtype to triton type, with reduction upcasts)r„  r�  r‚  s    r†   Útriton_acc_typer�  ä  s   € äÔ/°Ó6Ó7Ð7rˆ   c                óF   • U R                   S:*  =(       a    U R                  $ )Nr!   )r‹  Úis_floating_pointr‚  s    r†   Úlow_precision_fpr’  é  s   € Ø�>‰>˜QÑ×: 5×#:Ñ#:Ð:rˆ   c                óž   • [        U [        5      (       d  gU R                  n[        U[        R                  5      (       a  [	        U5      $ S$ rŒ  )rŠ   rT   rƒ  rœ   r’  )rÚ   rƒ  s     r†   Úlow_precision_fp_varr”  í  s<   € Ü�cœ;×'Ñ'Øà�I‰I€EÜ&0°¼¿¹×&DÑ&DÔ˜EÓ"ÐOÈ%ÐOrˆ   c                  óJ   ^ • \ rS rSr S         SU 4S jjjrS rSrU =r$ )ÚTritonCSEVariableiõ  c                óh   >• [         TU ]  XX4S9  [        5       U l        Uc   S5       eUc   S5       eg )N©rÉ   z!TritonCSEVariable must have dtypez!TritonCSEVariable must have shape)ÚsuperÚ__init__r   rû   )r  rÈ   Úboundsrƒ  rÉ   Ú	__class__s        €r†   rš  ÚTritonCSEVariable.__init__ö  sD   ø€ ô 	‰Ñ˜ uÐÑ:ä*4«,ˆŒØÑ ÐEÐ"EÓEÐ ØÑ ÐEÐ"EÓEÑ rˆ   c                óp  • U H°  n[        U[        5      (       a'  U R                  R                  UR                  5        M?  [        U[        R
                  5      (       d  M`  [        R                   H<  n[        XE5      (       d  M  U R                  R                  [        U    S3/5          M®     M²     g )Nr  )
rŠ   r–  rû   Úupdater‹   rô   r¸   rÐ   r   r   )r  rÈ   rÂ  ÚkwargsÚargrÝ   s         r†   Úupdate_on_argsÚ TritonCSEVariable.update_on_args  sƒ   € ÛˆCÜ˜#Ô0×1Ñ1Ø—‘×%Ñ% c§m¡mÖ4Ü˜C¤§¡×.Ó.ô *×5Ô5�DÜ% c×0Ó0ØŸ™×-Ñ-´*¸TÑ2BÐ1CÀ4Ð/HÐ.IÔJÚó 6ò rˆ   )rû   rš   )
rÈ   rÔ   r›  zValueRanges[Any]rƒ  útorch.dtyperÉ   rr   r¦   r§   )r›   r¨   r©   rª   rš  r¢  r®   Ú__classcell__©rœ  s   @r†   r–  r–  õ  sS   ø† ð !%ðFàðFð !ðFð ð	Fð
 ðFð 
÷Fð F÷ð rˆ   r–  c                 ó   • SSK Jn   U " 5       $ )Nr   rn   )Ú!torch._inductor.dtype_propagationro   rn   s    r†   Úget_dtype_handlerr©    s   € ÝLá%Ó'Ð'rˆ   c                ó:   ^ ^^• SS jmSU4S jjmSU UU4S jjnU$ )z‡
Codegen helper to upcast arguments to float32, depending on the config and dtype.
This decorates tl.math/libdevice codegen functions.
c                óÖ   • [         R                  R                  (       + =(       aD    [        U [        5      =(       a-    U R
                  [        R                  [        R                  4;   $ rš   )	r"   rµ   Úcodegen_upcast_to_fp32rŠ   rT   rƒ  rœ   Úfloat16Úbfloat16)rÚ   s    r†   Úneeds_upcastÚ*maybe_upcast_float32.<locals>.needs_upcast  sD   € ä—‘×4Ñ4Ô4÷ =Ü˜3¤Ó,÷=à—	‘	œeŸm™m¬U¯^©^Ð<Ñ<ð	
rˆ   c                ó2   >• T" U 5      (       a  SOSnU  U 3$ )Nú.to(tl.float32)r°   r‚   )rÚ   Úupcast_stringr¯  s     €r†   Úmaybe_upcast_argÚ.maybe_upcast_float32.<locals>.maybe_upcast_arg$  s$   ø€ Ù-9¸#×->Ñ->Ñ)ÀBˆØ��}�oÐ&Ð&rˆ   c                óL   >^ • [         R                  T T5        SUU UU4S jjnU$ )Nc                 óö  >• U  Vs/ s H  nT" U5      PM     nnUR                  5        VVs0 s H  u  pEUT" U5      _M     nnnT" U0 UD6nT=(       a7    [        U4S j[        R                  " XR	                  5       5       5       5      nU(       d  S O#[        [        5       TR                  5      " U 0 UD6n	U	[        R                  S 4;  n
U
(       a  U	b  S[        U	5       S3OSnU U 3$ s  snf s  snnf )Nc              3  ó4   >#   • U  H  nT" U5      v •  M     g 7frš   r‚   )rö   rÚ   r¯  s     €r†   r  ÚKmaybe_upcast_float32.<locals>.decorator.<locals>.wrapped.<locals>.<genexpr>3  s   øé € ð 6Ú-S c‘˜S×!Ð!Ò-Sùó   ƒú.to(r¡  r°   )Úitemsrz   Ú	itertoolsÚchainÚvaluesÚgetattrr©  r›   rœ   r�   rF   )rÂ  r   r¡  Úupcast_argsÚkeyÚvalÚupcast_kwargsrk  Úany_needs_upcastÚresult_dtypeÚneeds_downcastÚdowncast_stringr¡   r    r´  r¯  s               €€€€r†   ÚwrappedÚ8maybe_upcast_float32.<locals>.decorator.<locals>.wrapped,  s  ø€ á<@ÓAºD°SÑ+¨CÖ0¹DˆKÐAØHNÏÉÌÔWÊ¹H¸C˜SÑ"2°3Ó"7Ò7ÉˆMÑWñ ˜;Ð8¨-Ñ8ˆFØ-÷  ´#ô 6Ü-6¯_ª_¸TÇ=Á=Ã?Ô-Só6ó 3Ðö
 (ñ äÔ.Ó0°$·-±-Ô@À$ÐQÈ&ÑQð ð
 *´%·-±-ÀÐ1FÑFˆNö " lÑ&>ð ”{ <Ó0Ð1°Ñ3àð ð
 �X˜oÐ.Ð/Ð/ùò' BùÛWs
   †C0­C5r&  )r•   r£   )r    rÉ  r¡   r´  r¯  s   ` €€€r†   Ú	decoratorÚ'maybe_upcast_float32.<locals>.decorator(  s$   ù€ ä×&Ñ& t¨^Ô<÷	0ò 	0ð. ˆrˆ   r%  r&  )r    úCallable[..., Any]r¦   rÍ  r‚   )r¡   rË  r´  r¯  s   ` @@r†   Úmaybe_upcast_float32rÎ    s   ú€ ô
÷'÷ñ ð: Ðrˆ   c                  ó  • \ rS rSrSr\R                  " \R                  5      r\	  SR   SSS jj5       r
\	STS j5       r\	S 5       r\S 5       r\	\" 5       S	 5       5       r\	S
 5       r\	S 5       r\	\" 5       S 5       5       r\	\" 5       S 5       5       r\	\" 5       S 5       5       r\	\" 5       S 5       5       r\	S 5       r\	S 5       r\	S 5       r\	S 5       r\	S 5       r\	S\R:                  SSS.S j5       r\	\" 5       S 5       5       r\	\" 5       S 5       5       r \S 5       r!\	S 5       r"\	\" 5       S 5       5       r#\	\" 5       S 5       5       r$\	\" 5       S 5       5       r%\	\" 5       S 5       5       r&\	\" 5       S  5       5       r'\	\" 5       S! 5       5       r(\	\" 5       S" 5       5       r)\	\" 5       S# 5       5       r*\	\" 5       S$ 5       5       r+\	\" 5       S% 5       5       r,\	\" 5       S& 5       5       r-\	\" 5       S' 5       5       r.\	\" 5       S( 5       5       r/\	\" 5       S) 5       5       r0\	\" 5       S* 5       5       r1\	\" 5       S+ 5       5       r2\	\" 5       S, 5       5       r\	S- 5       r3\	\" 5       S. 5       5       r4\	S/ 5       r5\	S0 5       r6\	S1 5       r7\	S2 5       r8\	S3 5       r9\	S4 5       r:\	S5 5       r;\	S6 5       r<\	S7 5       r=\	S8 5       r>\	S9 5       r?\	S: 5       r@\	S; 5       rA\	S< 5       rB\	\" 5       S= 5       5       rC\	\" 5       S> 5       5       rD\	\" 5       S? 5       5       rE\	\" 5       S@ 5       5       rF\	\" 5       SA 5       5       rG\	SB 5       rH\	\" 5       SC 5       5       rI\	\" 5       SD 5       5       rJ\	\" 5       SE 5       5       rK\	\" SFSG9SH 5       5       rL\	\" SFSG9SI 5       5       rM\	\" 5       SJ 5       5       rN\	\" 5       SK 5       5       rO\	SL 5       rP\	SM 5       rQ\	\" 5       SN 5       5       rR\	SO 5       rS\	\" 5       SP 5       5       rTSQrUg)UÚTritonOverridesiH  zEMap element-wise ops to Triton e.g., ops.to_dtype(x,...) -> x.to(...)NTc                óx  •       SS jnUb=  [        U" X!5      [        R                  R                  5      [        R                  l        U[        R
                  :X  a  SU  S3$ U[        R                  :X  a  Ub  UR                  (       d  Uc  U  S3$ U(       a  [        U5      nO[        U5      nU  SU S3$ )Nc                ó   • X:X  a  g[         R                  [         R                  4nX;   a  X;   a  X:w  a   S5       eU [         R                  :X  d  U[         R                  :X  a  gU [         R                  :X  d  U[         R                  :X  a  gg)Nr   zCConversions between float8_e5m2 and float8_e4m3fn is not supported!r‰  r!   )rœ   Úfloat8_e4m3fnÚfloat8_e5m2)Ú	src_dtypeÚ	dst_dtypeÚ
fp8_dtypess      r†   Ú_get_min_elements_per_threadÚ>TritonOverrides.to_dtype.<locals>._get_min_elements_per_threadT  s”   € ð Ó%àô ×#Ñ#Ü×!Ñ!ðˆJð Ó'ØÓ+ØÓ*ðUð Uó	Uð ð
 œE×-Ñ-Ó-°¼e×>OÑ>OÓ1OØØœE×/Ñ/Ó/°9Ä×@SÑ@SÓ3SØàrˆ   r)  z != 0)z.to(tl.int16).to(tl.uint8)r»  r¡  )rÕ  r¤  rÖ  r¤  r¦   rî  )
r{   rL   rÅ   Úmin_elem_per_threadrœ   r¥   Úuint8r‘  r„  r‡  )rN  rƒ  rÕ  Úuse_compute_typesrØ  Ú	out_dtypes         r†   Úto_dtypeÚTritonOverrides.to_dtypeM  sÄ   € ð	Ø"ð	Ø/:ð	àô	ð6 Ñ ô ,/Ù,¨YÓ>Ü—‘×,Ñ,ó,ŒA�H‰HÔ(ð
 ”E—J‘JÓØ�q�c˜�=Ð Ø”e—k‘kÓ!ØÑ! i×&A×&AÀYÑEVð �SÐ2Ð3Ð3æÜ+¨EÓ2‰Iä)¨%Ó0ˆIà��D˜˜ 1Ð%Ð%rˆ   c                óô   • UR                   UR                   :X  d   eU R                  U:w  a  U  S[        U5       S3n U  S[        U5       S3n[        U5      U:w  a  U S[        [        U5      5       S3nU$ )Nr»  r¡  z, bitcast=True))r‹  rƒ  rF   rH   )rN  rƒ  rÕ  Úouts       r†   Úto_dtype_bitcastÚ TritonOverrides.to_dtype_bitcastŠ  sˆ   € à×!Ñ! U§^¡^Ó3Ð3Ð3ð �7‰7�iÓØ�#�Tœ+ iÓ0Ð1°Ð3ˆAà��4œ EÓ*Ð+¨?Ð;ˆÜ˜uÓ%¨Ó.Ø�E˜œkÔ*=¸eÓ*DÓEÐFÀaÐHˆCàˆ
rˆ   c           	     óì   • [         R                  R                  U5      n[        U" U 5      5      n[	        U5      nU S:  a(  UR
                  (       d  SUSS   3nSU SU SU SU S3	$ SU SU SU S3$ )Nr   ztl.r‰  útl.full(r   rã  r¡  )rœ   Ú_prims_commonÚdtype_to_typer`   r„  rŠ  )rª  rƒ  rÉ   Útype_Ú
triton_valrF   Útriton_signed_types          r†   Ú_shaped_constantÚ TritonOverrides._shaped_constant›  s’   € ä×#Ñ#×1Ñ1°%Ó8ˆÜ"¡5¨£<Ó0ˆ
Ü)¨%Ó0ˆð
 �1‹9˜UŸ_Ÿ_Ø#& {°1°2 Ð&7Ð!8ÐØ˜e˜W B z l°"Ð5GÐ4HÈÈkÈ]ÐZ[Ð\Ð\à˜e˜W B z l°"°[°MÀÐCÐCrˆ   c                ó"   • U R                  X/ S9$ )Nr˜  )rë  )rŸ   rª  rƒ  s      r†   ÚconstantÚTritonOverrides.constantª  s   € à×#Ñ# E¸Ð#Ð;Ð;rˆ   c                ó   • SU  S3$ )NrQ  r¡  r‚   ©rN  s    r†   ÚabsÚTritonOverrides.abs®  ó   € ð ˜a˜S Ð"Ð"rˆ   c                óÊ  • [        U SS 5      n[        USS 5      nU[        R                  :X  a=  U[        R                  :X  a)  [        R                  R
                  (       a
  SU  SU S3nO	SU  SU S3n[        U 5      (       d  [        U5      (       aN  [        5       R                  X5      nU[        R                  [        R                  4;   a  U S[        U5       S3nU$ )Nrƒ  ztriton.language.div_rn(r   r¡  r)  r  r»  )rÀ  rœ   r�   r"   Úeager_numericsÚdivision_roundingr”  r©  Útruedivr­  rF   )rN  ÚyÚx_dtypeÚy_dtyperá  rÝ  s         r†   rø  ÚTritonOverrides.truediv¶  sË   € ä˜!˜W dÓ+ˆÜ˜!˜W dÓ+ˆð ”u—}‘}Ó$Øœ5Ÿ=™=Ó(Ü×%Ñ%×7×7ð ,¨A¨3¨b°°°1Ð5‰Cà�a�S˜˜A˜3˜a�.ˆCä ×"Ñ"Ô&:¸1×&=Ñ&=Ü)Ó+×3Ñ3°AÓ9ˆIØœUŸ]™]¬E¯M©MÐ:Ó:Ø˜˜T¤+¨iÓ"8Ð!9¸Ð;�àˆ
rˆ   c                óô   • SU  SU S3n[        U 5      (       d  [        U5      (       aN  [        5       R                  X5      nU[        R                  [        R
                  4;   a  U S[        U5       S3nU$ )Nr)  rü  r¡  r»  )r”  r©  Úmodrœ   r­  r�   rF   )rN  rù  rá  rÝ  s       r†   rþ  ÚTritonOverrides.modÍ  sp   € à�!��C˜�s˜!ˆnˆÜ ×"Ñ"Ô&:¸1×&=Ñ&=Ü)Ó+×/Ñ/°Ó5ˆIØœUŸ]™]¬E¯M©MÐ:Ó:Ø˜˜T¤+¨iÓ"8Ð!9¸Ð;�Øˆ
rˆ   c                óD   • [         R                  (       a  SU  S3$ SU  S3$ )z 
When use_fast_math, use the ftz (flushing to zero) variant
of exponent computation.

Check https://github.com/triton-lang/triton/issues/5735 for
more details.
ztl_math.exp(r¡  zlibdevice.exp()r"   Úuse_fast_mathrñ  s    r†   r%  ÚTritonOverrides.expÖ  s+   € ô ××Ø! !  AÐ&Ð&à# A 3 aÐ(Ð(rˆ   c                ó   • SU  S3$ )Nzlibdevice.exp2(r¡  r‚   rñ  s    r†   Úexp2ÚTritonOverrides.exp2å  ó   € ð !   1Ð%Ð%rˆ   c                ó   • SU  S3$ )Nzlibdevice.expm1(r¡  r‚   rñ  s    r†   Úexpm1ÚTritonOverrides.expm1ê  ó   € ð " !  AÐ&Ð&rˆ   c                ó   • SU  S3$ )Nztl.sqrt_rn(r¡  r‚   rñ  s    r†   ÚsqrtÚTritonOverrides.sqrtï  ó   € ð ˜Q˜C˜qÐ!Ð!rˆ   c                ó  • [         R                  R                  nUS:X  a  gUS:X  a	  SU  SU  S3$ US:X  a  U  S3$ Uc:  [        R                  " [        R
                  " S	[        R                  5      U 5      $ [        S
U< 35      e)NÚcompile_errorzcompile error!Úruntime_errorz"triton_helpers.device_assert_then(z == 0, "injected assert fail", r¡  Úaccuracyz + 1r   z:unrecognized config triton.inject_relu_bug_TESTING_ONLY = )	r"   rµ   Úinject_relu_bug_TESTING_ONLYÚopsÚmaximumrî  rœ   rŒ  r`  )rN  Úbugs     r†   ÚreluÚTritonOverrides.reluô  s�   € ä�m‰m×8Ñ8ˆØ�/Ó!Ø#Ø�OÓ#ð 8¸°sÐ:YÐZ[ÐY\Ð\]Ð^Ð^Ø�JÓØ�S˜�:ÐØ‰[Ü—;’;œsŸ|š|¨A¬u¯{©{Ó;¸QÓ?Ð?ä ØLÈSÉGÐTóð rˆ   c                ód   • [         R                  R                  (       a	  SU  SU S3$ SU  SU S3$ )Nztl.minimum(r   ú, tl.PropagateNan.ALL)ztriton_helpers.minimum(r¡  ©rœ   rð  rñ  ©rD  rE  s     r†   ÚminimumÚTritonOverrides.minimum  ó:   € ä�=‰=××Ø    2 a SÐ(>Ð?Ð?à,¨Q¨C¨r°!°°AÐ6Ð6rˆ   c                ód   • [         R                  R                  (       a	  SU  SU S3$ SU  SU S3$ )Nztl.maximum(r   r  ztriton_helpers.maximum(r¡  r  r  s     r†   r  ÚTritonOverrides.maximum  r  rˆ   c                ó   • SU  SU SU S3$ )Nr1  r   r¡  r‚   )rD  rE  r3  s      r†   ÚwhereÚTritonOverrides.where  s   € à˜1˜#˜R ˜s " Q C qÐ)Ð)rˆ   c           	     ó¸  ^^	^
^• [         R                  R                  (       d   eXp2S nS nU" U5      (       a  U" U 5      n U" U5      (       a  U" U5      n      SUU	U
U4S jjn[        [         R                  R	                  5       5      S:¼  d   S5       e[        [        R                  [        R                     5      m	[        [        R                  [        R                     5      m
[        [        R                  [        R                     5      m[        [        R                  [        R                     5      m[         R                  R                  R                  [         R                  R                  U" U [!        U R"                  5      T
T/5      U R$                  T
T4S9n [         R                  R                  R                  [         R                  R                  U" U[!        UR"                  5      TT	/5      UR$                  TT	4S9n[&        R(                  R*                  R,                  R.                  S:X  a  SnOSnS	U  S
U SU S3$ )at  
Triton code generation for lowering ops.dot to tl.dot.

The logic is as follows:

1. Downcasting for performance
   If the data was previously upcasted to fp32, we downcast back to the
   original dtype (e.g., fp16 or bf16) for better performance. While
   surrounding operations may run in fp32, matmul itself is executed at the
   original precision to optimize throughput.

2. Handling non-constant reduction masks
   If the reduction mask is not constant and there was any operation between
   tl.load and tl.dot, we zero out regions outside the mask using
   tl.where(r0_mask, val, 0).
   This ensures that values outside the mask do not contribute to the dot
   product, preventing incorrect results.

3. Shape alignment for tl.dot
   We massage shapes to match the tl.dot requirement of (Y, R) x (R, X).
   Current codegen eagerly broadcasts tl.arange to create unique axes. We
   reshape, transpose, or broadcast to align with the (Y, R) x (R, X) shape.
   We avoid using 3D dot ((Z, Y, R) x (Z, R, X)) because 3D tl.dot has
   poor performance. During batched matmul (bmm), we keep ZBLOCK=1 and call
   the 2D dot kernel instead.
c                ó¤  • [        [        [        U R                  5      5      (       d  g[        R
                  R                  S   nUR                  (       d   e[        R
                  R                  U5      (       a  g[        R
                  R                  R                  R                  5        H  u  p#X0:X  d  M  SU;   d  M  SU;   d  M    g   g)NFrK  ztl.loadz	other=0.0T)rz   r"  rB   rû   rL   rÅ   re  r\  Ú_has_constant_maskrÆ   Ú_cacher¼  )rÚ   Úreduction_rangeÚkÚvs       r†   Úis_where_neededÚ,TritonOverrides.dot.<locals>.is_where_needed7  s˜   € ä”sÔ.°·±Ó>×?Ñ?ØäŸh™h×2Ñ2°2Ñ6ˆOØ"×/×/Ð/Ð/ô �x‰x×*Ñ*¨?×;Ñ;Øô
 Ÿ™Ÿ™×+Ñ+×1Ñ1Ö3‘�Ø•8 	¨Q¥°;À!Õ3CÙ ñ 4ð rˆ   c                óü  • [         R                  R                  SU R                  5      n[        R
                  R                   Vs/ s H%  nUR                  (       d  M  UR                   S3PM'     nn[        U5      S:X  d   S5       e[        R                  US   X5      n[        R
                  R                  R                  [        R
                  R                  X@R                  U R                  S9$ s  snf )NÚdotr  rN   z'don't tile reduction when native matmulr   ©rƒ  rÉ   )r#   Ú	ReductionÚdefault_valuerƒ  rL   rÅ   re  r\  rÓ   r�   ÚTritonKernelOverridesr#  rÆ   ÚgenerateÚcomputerÉ   )rÚ   Údefaultrá   Úreduction_maskÚ	where_vars        r†   Ú
where_condÚ'TritonOverrides.dot.<locals>.where_condL  sÎ   € Ü—l‘l×0Ñ0°¸¿	¹	ÓBˆGô ŸH™H×0Ò0óâ0�DØ×$Õ$ó %�4—;‘;�-˜tÓ$Ù0ð ð ô �~Ó&¨!Ó+ÐVÐ-VÓVÐ+ä-×3Ñ3°NÀ1Ñ4EÀsÓTˆIÜ—8‘8—<‘<×(Ñ(Ü—‘× Ñ  )·9±9ÀCÇIÁIð )ð ð ùòs   ÁC9Á C9c                óÀ  >• T	U;   a  U Vs/ s H  o3T	:X  a  SOUPM     nnUTT/:X  aO  TU;  d   S5       eSS/nTU;   a  TUS'   TU;   a  TUS'   [        XU5      n UTT/:g  nU(       a  SU  ST ST S3n U $ UTT/:X  aU  TU;  d   S	5       eSS/nTU;   a  TUS'   TU;   a  TUS'   [        XU5      n S
U  S3n UTT/:g  nU(       a  SU  ST ST S3n U $ [        es  snf )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
rÃ   z&left tl.dot operand cannot depend on xr   rN   rŸ  z, (r   r@  z'right tl.dot operand cannot depend on yr¢  r¡  )r¥  ÚNotImplementedError)
rª  r«  r3  rg  Úshape_2dÚbroadcast_neededÚRBLOCKrñ   rò   ró   s
         €€€€r†   Ú#reshape_transpose_broadcast_for_dotÚ@TritonOverrides.dot.<locals>.reshape_transpose_broadcast_for_dotl  sw  ø€ ð6 ˜Ó&ÙJWÓ XÊ-À3¨v«¡¸3Ò!>É-�Ð Xà˜v vÐ.Ó.Ø ]Ó2ð Ø<óÐ2ð   ˜:�Ø˜]Ó*Ø"(�H˜Q‘KØ˜]Ó*Ø"(�H˜Q‘Kô ' u¸XÓF�ð $,°¸Ð/?Ñ#?Ð Þ#Ø.¨u¨g°S¸¸ÀÀ6À(È"ÐM�Eð4 ˆLð1  ¨Ð 0Ó0Ø ]Ó2ð Ø=óÐ2ð   ˜:�Ø˜]Ó*Ø"(�H˜Q‘KØ˜]Ó*Ø"(�H˜Q‘Kô ' u¸XÓF�ð $ E 7¨!Ð,�ð $,°¸Ð/?Ñ#?Ð Þ#Ø.¨u¨g°S¸¸ÀÀ6À(È"ÐM�Eð ˆLô *Ð)ùòW !Ys   ŒCr   ztl.dot can only do mm and bmmr0  Útf32Úieeeztl.dot(r   z, input_precision="ú"))r«  r/  r3  r/  r¦   rÔ   )rL   rÅ   Úis_native_matmulr�   Údense_size_listrÔ   r¸   ræ   r   rñ   rò   ró   rî   rÆ   r4  r5  ra  rÉ   rƒ  rœ   ÚbackendsÚcudaÚmatmulÚfp32_precision)rD  rE  Úorig_aÚorig_br,  r9  r@  Úinput_precisionr?  rñ   rò   ró   s           @@@@r†   r/  ÚTritonOverrides.dot  sö  û€ ô8 �x‰x×(×(Ð(Ð(Ø�ò	ò*	ñ6 ˜6×"Ñ"Ù˜1“ˆAÙ˜6×"Ñ"Ù˜1“ˆAðI	à/ðI	ð .ðI	ð ÷	I	ò I	ôV ”1—8‘8×+Ñ+Ó-Ó.°!Ó3ÐTÐ5TÓTÐ3ä”]×.Ñ.¬t¯{©{Ñ;Ó<ˆÜ”]×.Ñ.¬t¯{©{Ñ;Ó<ˆÜ”]×.Ñ.¬t¯{©{Ñ;Ó<ˆÜ”]×.Ñ.¬t¯}©}Ñ=Ó>ˆä�H‰H�L‰L×!Ñ!Ü�H‰H×ÑÙ/°´4¸¿¹³=À6È6ÐBRÓSØ—'‘'Ø˜6Ð"ð	 "ð 
ˆô �H‰H�L‰L×!Ñ!Ü�H‰H×ÑÙ/°´4¸¿¹³=À6È6ÐBRÓSØ—'‘'Ø˜6Ð"ð	 "ð 
ˆô �>‰>×Ñ×%Ñ%×4Ñ4¸Ó>Ø$‰Oà$ˆOà˜˜˜2˜a˜SÐ 3°OÐ3DÀBÐGÐGrˆ   rN   )Úconstraintsrƒ  Úis_pureÚpackc                óþ   • [        U5      nSR                  U Vs/ s H  n[        U5      PM     sn5      nUc&  SR                  S/U V	s/ s H  n	SPM     sn	-   5      nSU  SU SU SU SU S	U S
3$ s  snf s  sn	f )Nr   z=rr  ztl.inline_asm_elementwise('z', 'z', [z	], dtype=z
, is_pure=z, pack=r¡  )r„  r   rÔ   )
ÚasmrO  rƒ  rP  rQ  ÚinputsrF   ÚiÚ
input_refsÚ_s
             r†   Úinline_asm_elementwiseÚ&TritonOverrides.inline_asm_elementwiseÓ  s®   € ô *¨%Ó0ˆØ—Y‘Y±Ó7²¨1¤ A¦±Ñ7Ó8ˆ
ØÑØŸ)™) T F¹6Ó-Bº6°a«c¹6Ñ-BÑ$BÓCˆKØ,¨S¨E°°k°]À$ÀzÀlÐR[Ð\gÐ[hÐhrÐszÐr{ð  |Cð  DHð  CIð  IJð  Kð  	Kùò  8ùâ-Bs   ›A5ÁA:
c                ó   • SU  S3$ )Nztl_math.cos(r¡  r‚   rñ  s    r†   ÚcosÚTritonOverrides.cosÝ  rô  rˆ   c                ó   • SU  S3$ )Nztl_math.sin(r¡  r‚   rñ  s    r†   ÚsinÚTritonOverrides.sinâ  rô  rˆ   c                ó   • [        S5      e)Nz/ops.index_expr not implemented outside a kernel©r<  )rŸ   r’   rƒ  s      r†   Ú
index_exprÚTritonOverrides.index_exprç  s   € ä!Ð"SÓTÐTrˆ   c                ó   • [        S5      e)Nz+ops.masked not implemented outside a kernelra  )r  ÚbodyÚothers      r†   ÚmaskedÚTritonOverrides.maskedë  s   € ä!Ð"OÓPÐPrˆ   c                ó   • SU  S3$ )Nzlibdevice.lgamma(r¡  r‚   rñ  s    r†   ÚlgammaÚTritonOverrides.lgammaï  ó   € ð # 1 # QÐ'Ð'rˆ   c                ó   • SU  S3$ )Nzlibdevice.erf(r¡  r‚   rñ  s    r†   ÚerfÚTritonOverrides.erfô  ó   € ð   ˜s !Ð$Ð$rˆ   c                ó   • SU  S3$ )Nzlibdevice.cosh(r¡  r‚   rñ  s    r†   ÚcoshÚTritonOverrides.coshù  r  rˆ   c                ó   • SU  S3$ )Nzlibdevice.sinh(r¡  r‚   rñ  s    r†   ÚsinhÚTritonOverrides.sinhþ  r  rˆ   c                ó   • SU  S3$ )Nzlibdevice.acos(r¡  r‚   rñ  s    r†   ÚacosÚTritonOverrides.acos  r  rˆ   c                ó   • SU  S3$ )Nzlibdevice.acosh(r¡  r‚   rñ  s    r†   ÚacoshÚTritonOverrides.acosh  r
  rˆ   c                ó   • SU  S3$ )Nzlibdevice.asin(r¡  r‚   rñ  s    r†   ÚasinÚTritonOverrides.asin  r  rˆ   c                ó   • SU  S3$ )Nzlibdevice.asinh(r¡  r‚   rñ  s    r†   ÚasinhÚTritonOverrides.asinh  r
  rˆ   c                ó   • SU  SU S3$ )Nzlibdevice.atan2(r   r¡  r‚   ©rN  rù  s     r†   Úatan2ÚTritonOverrides.atan2  ó   € ð " !  B q c¨Ð+Ð+rˆ   c                ó   • SU  S3$ )Nzlibdevice.atan(r¡  r‚   rñ  s    r†   ÚatanÚTritonOverrides.atan  r  rˆ   c                ó   • SU  S3$ )Nzlibdevice.atanh(r¡  r‚   rñ  s    r†   ÚatanhÚTritonOverrides.atanh!  r
  rˆ   c                ó   • SU  SU S3$ )Nzlibdevice.copysign(r   r¡  r‚   r„  s     r†   ÚcopysignÚTritonOverrides.copysign&  s   € ð % Q C r¨!¨¨AÐ.Ð.rˆ   c                ó   • SU  S3$ )Nzlibdevice.erfc(r¡  r‚   rñ  s    r†   ÚerfcÚTritonOverrides.erfc+  r  rˆ   c                ó   • SU  S3$ )Nzlibdevice.erfinv(r¡  r‚   rñ  s    r†   ÚerfinvÚTritonOverrides.erfinv0  rl  rˆ   c                ó   • SU  SU S3$ )Nzlibdevice.hypot(r   r¡  r‚   r„  s     r†   ÚhypotÚTritonOverrides.hypot5  r‡  rˆ   c                ó   • SU  S3$ )Nzlibdevice.log10(r¡  r‚   rñ  s    r†   Úlog10ÚTritonOverrides.log10:  r
  rˆ   c                ó   • SU  S3$ )Nzlibdevice.log2(r¡  r‚   rñ  s    r†   Úlog2ÚTritonOverrides.log2?  r  rˆ   c                ó   • SU  SU S3$ )Nzlibdevice.ldexp(r   z.to(tl.int32))r‚   )rN  Úns     r†   ÚldexpÚTritonOverrides.ldexpD  s   € à! !  B q c¨Ð8Ð8rˆ   c                ó   • SU  SU S3$ )Nzlibdevice.nextafter(r   r¡  r‚   r„  s     r†   Ú	nextafterÚTritonOverrides.nextafterH  s   € ð & a S¨¨1¨#¨QÐ/Ð/rˆ   c                ó   • U  SU 3$ ©Nr  r‚   r  s     r†   Úlogical_andÚTritonOverrides.logical_andM  ó   € à��C˜�sˆ|Ðrˆ   c                ó   • U  S3$ )Nz == 0r‚   ©rD  s    r†   Úlogical_notÚTritonOverrides.logical_notQ  s   € à��Eˆ{Ðrˆ   c                ó   • U  SU 3$ ©Nz | r‚   r  s     r†   Ú
logical_orÚTritonOverrides.logical_orU  r«  rˆ   c                ó   • SU  SU S3$ )Nr)  ú ^ r¡  r‚   r  s     r†   Úlogical_xorÚTritonOverrides.logical_xorY  s   € à�1�#�S˜˜˜1ˆ~Ðrˆ   c                ó   • U  SU 3$ r¨  r‚   r  s     r†   Úbitwise_andÚTritonOverrides.bitwise_and]  r«  rˆ   c                ó   • SU  3$ )NÚ~r‚   r­  s    r†   Úbitwise_notÚTritonOverrides.bitwise_nota  s   € à�1�#ˆwˆrˆ   c                ó   • U  SU 3$ r±  r‚   r  s     r†   Ú
bitwise_orÚTritonOverrides.bitwise_ore  r«  rˆ   c                ó   • U  SU 3$ )Nrµ  r‚   r  s     r†   Úbitwise_xorÚTritonOverrides.bitwise_xori  r«  rˆ   c                ó   • U  SU 3$ )Nz << r‚   r  s     r†   Úbitwise_left_shiftÚ"TritonOverrides.bitwise_left_shiftm  ó   € à��D˜˜ˆ}Ðrˆ   c                ó   • U  SU 3$ )Nz >> r‚   r  s     r†   Úbitwise_right_shiftÚ#TritonOverrides.bitwise_right_shiftq  rÈ  rˆ   c                ó    • SU S3nSU  SU S3$ )Nr)  ú).to(tl.uint32)ztl.rand(r   r¡  r‚   ©Úseedrº   s     r†   ÚrandÚTritonOverrides.randu  s%   € à�V�H˜OÐ,ˆØ˜$˜˜r & ¨Ð+Ð+rˆ   c                ó    • SU S3nSU  SU S3$ )Nr)  rÍ  z	tl.randn(r   r¡  r‚   rÎ  s     r†   ÚrandnÚTritonOverrides.randnz  s%   € à�V�H˜OÐ,ˆØ˜4˜&  6 (¨!Ð,Ð,rˆ   c           	     ó,   • SU S3nSU  SU SU SU S3	$ )Nr)  rÍ  ztriton_helpers.randint64(r   r¡  r‚   )rÏ  rº   ÚlowÚhighs       r†   Ú	randint64ÚTritonOverrides.randint64  s1   € à�V�H˜OÐ,ˆØ*¨4¨&°°6°(¸"¸S¸EÀÀDÀ6ÈÐKÐKrˆ   c                ó   • [        S5      e)Nz.ops.load_seed not implemented outside a kernelra  )rÈ   rº   s     r†   Ú	load_seedÚTritonOverrides.load_seed„  s   € ä!Ð"RÓSÐSrˆ   c                óX   • [         R                  R                  (       a  SU  S3$ SU  S3$ )Nz	tl.rsqrt(r¡  zlibdevice.rsqrt(r  rñ  s    r†   ÚrsqrtÚTritonOverrides.rsqrtˆ  s/   € ô �=‰=××Ø˜q˜c Ð#Ð#à% a S¨Ð*Ð*rˆ   c                ó   • SU  S3$ )Nzlibdevice.log1p(r¡  r‚   rñ  s    r†   Úlog1pÚTritonOverrides.log1p�  r
  rˆ   c                ó   • SU  S3$ )Nzlibdevice.tan(r¡  r‚   rñ  s    r†   ÚtanÚTritonOverrides.tan•  rp  rˆ   c                ó€  • [         R                  R                  R                  R	                  U 5      nU(       a  [        US5      (       a  UR                  nOS n[        R                  (       aJ  [        R                  R                  (       a+  [        5       S:”  a  U[        R                  :w  a	  Ub  SU  S3$ SU  S3$ )Nrƒ  )r   é   zlibdevice.fast_tanhf(r¡  zlibdevice.tanh()rL   rÅ   rÆ   rÇ   Úgetr³   rƒ  r"   r  rœ   rð  rñ  r   rž   )rN  rÛ   rƒ  s      r†   ÚtanhÚTritonOverrides.tanhš  s’   € ô —(‘(—,‘,×*Ñ*×.Ñ.¨qÓ1ˆÞ”w˜w¨×0Ñ0Ø—M‘M‰EàˆEä× × Ü—‘×!×!Ü"Ó$ vÓ-ØœŸ™Ó&ØÑ!ð +¨1¨#¨QÐ/Ð/à$ Q C qÐ)Ð)rˆ   c                ó   • SU  S3$ )Nztl.sigmoid(r¡  r‚   rñ  s    r†   ÚsigmoidÚTritonOverrides.sigmoid¯  r  rˆ   c                ó   • SU  SU  SU  S3$ )Nz(libdevice.signbit(z) != 0) if (z).dtype is tl.float32 else z < 0r‚   rñ  s    r†   ÚsignbitÚTritonOverrides.signbit´  s#   € ð " !  L°°Ð3NÈqÈcÐQUÐVð	
rˆ   c                ó   • SU  SU S3$ )Nzlibdevice.fmod(r   r¡  r‚   r  s     r†   ÚfmodÚTritonOverrides.fmod»  s   € ð !   2 a S¨Ð*Ð*rˆ   c                ó   • SU  SU S3$ )Nr*  r   r¡  r‚   r  s     r†   ÚpowÚTritonOverrides.powÀ  s   € ð   ˜s " Q C qÐ)Ð)rˆ   c                ó   • SU  S3$ )Nztl_math.log(r¡  r‚   rñ  s    r†   ÚlogÚTritonOverrides.logÅ  rô  rˆ   F)r¡   c                ó   • SU  S3$ )Nzlibdevice.isinf(ú).to(tl.int1)r‚   rñ  s    r†   ÚisinfÚTritonOverrides.isinfÊ  ó   € ð " !  MÐ2Ð2rˆ   c                ó   • SU  S3$ )Nzlibdevice.isnan(rû  r‚   rñ  s    r†   ÚisnanÚTritonOverrides.isnanÏ  rþ  rˆ   c                ó   • SU  S3$ )Nzlibdevice.nearbyint(r¡  r‚   rñ  s    r†   ÚroundÚTritonOverrides.roundÔ  s   € ð & a S¨Ð*Ð*rˆ   c                ó   • SU  S3$ )Nr  r¡  r‚   rñ  s    r†   ÚfloorÚTritonOverrides.floorÙ  r
  rˆ   c                óH   • U  SU 3nU  SU 3nSU  SU SU SU SU SU S	3$ )
Nr  rü  z
tl.where((z
 < 0) != (z < 0), tl.where(z != 0, z - 1, ú), r¡  r‚   )rD  rE  r   Úrems       r†   ÚfloordivÚTritonOverrides.floordivÞ  sV   € ð
 ��D˜˜ˆ}ˆØ��3�q�cˆlˆØ˜A˜3˜j¨¨Ð+;¸C¸5ÀÈÀvÈVÐTXÐSYÐY\Ð]aÐ\bÐbcÐdÐdrˆ   c                ól  • [         R                  " S[        R                  5      n[         R                  " [         R
                  " X5      [        R                  5      n[         R                  " [         R
                  " X5      [        R                  5      n[         R                  " X#5      nU SU  S3$ )Nr   r»  ú.dtype))r  rî  rœ   rŒ  rÞ  Últr†  Úsub)rN  ÚzÚleftÚrightr  s        r†   ÚsignÚTritonOverrides.signç  so   € ä�LŠL˜œEŸK™KÓ(ˆÜ�|Š|œSŸVšV A›\¬E¯J©JÓ7ˆÜ—’œcŸfšf Q›l¬U¯Z©ZÓ8ˆÜ�gŠg�dÓ"ˆØ��d˜1˜#˜WÐ%Ð%rˆ   c                ó   • SU  S3$ )Nrâ  r¡  r‚   rñ  s    r†   ÚtruncÚTritonOverrides.truncï  r
  rˆ   c                ó   • U  SU 3$ )Nr  r‚   r  s     r†   ÚtruncdivÚTritonOverrides.truncdivô  s   € ð ��D˜˜ˆ}Ðrˆ   c                ó   • SU  S3$ )Nr  r¡  r‚   rñ  s    r†   ÚceilÚTritonOverrides.ceilú  r  rˆ   r‚   ©NT)rƒ  r¤  rÕ  úOptional[torch.dtype])rƒ  r¤  rÕ  r¤  )Vr›   r¨   r©   rª   r«   Úmathrž  ÚeÚ_LOG_2_EÚstaticmethodrÞ  râ  rë  r­   rî  rÎ  rò  rø  rþ  r%  r  r  r  r  r  r  r#  r/  rœ   r�   rX  r[  r^  rb  rg  rj  rn  rr  ru  rx  r{  r~  r�  r…  r‰  rŒ  r�  r’  r•  r˜  r›  r¢  r¥  r©  r®  r²  r¶  r¹  r½  rÀ  rÃ  rÆ  rÊ  rÐ  rÓ  rØ  rÛ  rÞ  rá  rä  ré  rì  rï  rò  rõ  rø  rü  r   r  r  r  r  r  r  r  r®   r‚   rˆ   r†   rÐ  rÐ  H  s¢  † ÙOà�yŠy˜Ÿ™Ó €Hàð ,0Øð	:&àð:&ð )ô:&ó ð:&ðx óó ðð  ñDó ðDð ñ<ó ð<ð ÙÓñ#ó ó ð#ð ñó ðð, ñó ðð ÙÓñ)ó ó ð)ð ÙÓñ&ó ó ð&ð ÙÓñ'ó ó ð'ð ÙÓñ"ó ó ð"ð ñó ðð" ñ7ó ð7ð ñ7ó ð7ð ñ*ó ð*ð ñxHó ðxHðt à"&¨e¯m©mÀTÐPQôKó ðKð ÙÓñ#ó ó ð#ð ÙÓñ#ó ó ð#ð ñUó ðUð ñQó ðQð ÙÓñ(ó ó ð(ð ÙÓñ%ó ó ð%ð ÙÓñ&ó ó ð&ð ÙÓñ&ó ó ð&ð ÙÓñ&ó ó ð&ð ÙÓñ'ó ó ð'ð ÙÓñ&ó ó ð&ð ÙÓñ'ó ó ð'ð ÙÓñ,ó ó ð,ð ÙÓñ&ó ó ð&ð ÙÓñ'ó ó ð'ð ÙÓñ/ó ó ð/ð ÙÓñ&ó ó ð&ð ÙÓñ(ó ó ð(ð ÙÓñ,ó ó ð,ð ÙÓñ'ó ó ð'ð ÙÓñ&ó ó ð&ð ñ9ó ð9ð ÙÓñ0ó ó ð0ð ñó ðð ñó ðð ñó ðð ñó ðð ñó ðð ñó ðð ñó ðð ñó ðð ñó ðð ñó ðð ñ,ó ð,ð ñ-ó ð-ð ñLó ðLð ñTó ðTð ÙÓñ+ó ó ð+ð ÙÓñ'ó ó ð'ð ÙÓñ%ó ó ð%ð ÙÓñ*ó ó ð*ð& ÙÓñ"ó ó ð"ð ñ
ó ð
ð ÙÓñ+ó ó ð+ð ÙÓñ*ó ó ð*ð ÙÓñ#ó ó ð#ð Ù¨Ñ/ñ3ó 0ó ð3ð Ù¨Ñ/ñ3ó 0ó ð3ð ÙÓñ+ó ó ð+ð ÙÓñ'ó ó ð'ð ñeó ðeð ñ&ó ð&ð ÙÓñ'ó ó ð'ð ñó ðð
 ÙÓñ&ó ó ó&rˆ   rÐ  rµ   c                  óÒ   ^ • \ rS rSrSrU 4S jr\\R                  S 5       5       r	\S 5       r
\S 5       r\S 5       r\S 5       r\S	 5       r\          SS
 j5       rSrU =r$ )r3  i  a  Map element-wise ops to Triton within a TritonKernel

Unlike TritonOverrides, these assume the code is going to be inserted into
the body of the main triton kernel and so it may use indexing and mask
variables which are assumed to already be defined in the current scope.
c                óF   >• [         TU ]  " U0 UD6  U R                  5         g rš   )r™  rš  Ú_setup_libdevice_routing)r  rÂ  r   rœ  s      €r†   rš  ÚTritonKernelOverrides.__init__  s#   ø€ Ü‰Ò˜$Ð) &Ò)ð 	×%Ñ%Õ'rˆ   c                óª  ^• SSK Jm  [        R                  R                  R
                   H¤  n[        X5      (       d   e[        X5      nU4S jnUS:X  aE  [        TS5      (       d   e[        R                  " X2US9nXl
        [        X[        U5      5        Mq  S n[        R                  " XRUS9nXl
        [        X[        U5      5        M¦     g)z<Set up routing to libdevice implementations for fp64 inputs.r   )ÚOpDecompositionsc                óˆ   >• U R                   [        R                  :w  a  U" U 5      $ [        TU5      " U 5      R                  $ rš   )rƒ  rœ   rž   rÀ  rª  )rN  Ú_original_implÚ_fn_namer*  s      €r†   Údecomposition_routerÚLTritonKernelOverrides._setup_libdevice_routing.<locals>.decomposition_router  s7   ø€ Ø—7‘7œeŸm™mÓ+Ù)¨!Ó,Ð,ä"Ð#3°XÔ>¸qÓA×GÑGÐGrˆ   rì  )r,  r-  c                ó`   • U R                   [        R                  :X  a	  SU SU  S3$ U" U 5      $ )Nz
libdevice.r)  r¡  )rƒ  rœ   rž   )rN  r,  r-  s      r†   Údtype_routerÚDTritonKernelOverrides._setup_libdevice_routing.<locals>.dtype_router,  s2   € Ø—7‘7œeŸm™mÓ+Ø'¨ z°°1°#°QÐ7Ð7á)¨!Ó,Ð,rˆ   N)Útorch._inductor.codegen.commonr*  rœ   Ú	_inductorr%   Úop_requires_libdevice_fp64r³   rÀ  Ú	functoolsÚpartialr›   Úsetattrr$  )rŸ   Úfn_nameÚoriginal_implr.  Úfnr1  r*  s         @r†   r'  Ú.TritonKernelOverrides._setup_libdevice_routing  sÄ   ø€ õ
 	Dä—‘×,Ñ,×GÔGˆGÜ˜3×(Ñ(Ð(Ð(Ü# CÓ1ˆMõHð ˜)Ó#ÜÐ/°×;Ñ;Ð;Ð;Ü×&Ò&Ø(ÐQXñ�ð &”Ü˜¤l°2Ó&6Ô7Ùò-ô ×"Ò"ØÀWñˆBð "ŒKÜ�C¤,¨rÓ"2Ö3ò; Hrˆ   c                ój   • [         R                  R                  5       nS/U-  nU R                  XUS9$ )NrN   r˜  )rL   rÅ   rÑ   rë  )rŸ   rª  rƒ  rà   rÉ   s        r†   rî  ÚTritonKernelOverrides.constant8  s7   € ô
 �x‰x×*Ñ*Ó,ˆØ��d‘
ˆØ×#Ñ# E¸Ð#Ð>Ð>rˆ   c                óâ  • [         R                  R                  USS S9n[        U[        5      (       d   eUR
                  (       a  UR
                  nO[        R                  UR                  5      n[         R                  R                  5       nU[        R                  [        R                  4;  a  UOUn[        R                  R                  n S[        R                  l        [         R                  R                   R#                  [         R                  R$                  UR&                  [)        U5      UUS9nU[        R                  l        U[        R                  [        R                  4;  ae  [         R                  R                   R#                  [         R                  R$                  U R+                  Xr5      [-        U5      UR.                  S9nOæUnUR0                   Ht  n[3        U[4        R6                  5      (       d  M$  [        R8                  " U[         R                  R                   R:                  UR<                     R>                  5      nMv     X%:w  a[  [         R                  R                   R#                  [         R                  R$                  U R+                  Xu5      UUR.                  S9nUR@                  Ul         U$ ! U[        R                  l        f = f)NF©Ú	block_ptrÚtma_compatibility_checker©r›  rƒ  rÉ   r0  )!rL   rÅ   ÚindexingrŠ   rø   rÿ   r¸   rã   rþ   Úget_index_dtype_as_torch_dtyperœ   rŒ  Úint64r"   Útest_configsÚruntime_triton_dtype_assertrÆ   r4  r5  rù   r=   rÞ  rH   rÉ   r�   r   r   rÄ   Úpromote_typesrÇ   rÈ   rƒ  rû   )	rŸ   r’   rƒ  rD  rÉ   ræ  ÚorigrÚ   Ú	index_vars	            r†   rb  Ú TritonKernelOverrides.index_exprA  s,  € ä—8‘8×$Ñ$Ø˜E¸Tð %ð 
ˆô ˜(¤O×4Ñ4Ð4Ð4ð × × Ø×)Ñ)‰Eä!×1Ñ1°(·.±.ÓAˆEô —h‘h×=Ñ=Ó?ˆØ¤u§{¡{´E·K±KÐ&@Ó@‘Àkˆô ×"Ñ"×>Ñ>ˆð
	CØ>CŒF×ÑÔ;Ü—(‘(—,‘,×'Ñ'Ü—‘× Ñ Ø×"Ñ"Ü,¨TÓ2ØØð (ð ˆCð ?CŒF×ÑÔ;àœŸ™¤e§k¡kÐ2Ó2Ü—(‘(—,‘,×'Ñ'Ü—‘× Ñ Ø—‘˜SÓ(Ü)¨%Ó0Ø—i‘ið	 (ð ‰Cð  ˆEØ!×.Ô.�	Ü! )¬T¯X©X×6Ó6Ü!×/Ò/ØœqŸx™xŸ|™|×7Ñ7¸	¿¹ÑG×MÑMó’Eñ /ð Ó#Ü—h‘h—l‘l×+Ñ+Ü—H‘H×$Ñ$Ø—L‘L Ó2Ø%ØŸ)™)ð	 ,ð �ð !×*Ñ*ˆŒØˆ
øð= ?CŒF×ÑÕ;ús   ÃA+K ËK.c           
     ó<  • U bx  [         R                  R                  b]  [        R                  R
                  R                  [        R                  R                  U  S3[         R                  U R                  S9n UR                  R                  SS9nU(       d   S5       eSnU HH  nUR                   H5  nUR                  S:w  d  [        UR                  S   5      (       d  M2  S	n  MF     MJ     U(       a  S OUn[        R                  R                  XS
9 nU" 5       n	S S S 5        U(       a¼  W	R                   R"                  (       a  [        U5      n[        R                  R
                  R                  [        R                  R                  SU	 S[%        U5       SU	 S3[&        R(                  " U5      U	R*                  U	R                  S9n[,        R.                  " WX’5      n
OW	n
U
R0                  R3                  W5        U
$ ! , (       d  f       Nð= f)Nú.to(tl.int1)r0  Úoutput)Úopz)graph for body does not contain an outputFÚloadrN   T©rª  rå  z.shape, r   r  rC  )rœ   rð  rñ  rL   rÅ   rÆ   r4  r5  r¥   rÉ   rP  Ú
find_nodesrÂ  Útargetrj   Ú
mask_loadsr›  Úis_boolr`   r    Úwraprƒ  r  r#  rû   Údiscard)r  re  rf  ÚnodesÚ
need_whereÚnoder¡  rª  Únew_maskrk  rò  s              r†   rg  ÚTritonKernelOverrides.masked  sª  € àÑ¤§¡× 1Ñ 1Ñ =Ü—8‘8—<‘<×(Ñ(Ü—‘× Ñ Ø�&˜Ð%Ü—j‘jØ—j‘jð	 )ð ˆDð —
‘
×%Ñ%¨Ð%Ð2ˆÞÐAÐAÓAˆuàˆ
ó ˆDØ—y”y�Ø—:‘: Ó'Ô+CÀCÇHÁHÈQÁK×+PÓ+PØ!%�JÚó !ñ ö #‘¨ˆä�X‰X× Ñ  Ð Ñ3°xÙ“VˆF÷ 4ö à�}‰}×$×$Ü˜U›�ä—H‘H—L‘L×)Ñ)Ü—‘× Ñ Ø˜6˜( (¬=¸Ó+?Ð*@ÀÀ6À(È'ÐRÜ"×'Ò'¨Ó.Ø—l‘lØ—l‘lð *ð ˆEô —)’)˜H fÓ4‰CàˆCà�‰×Ñ˜hÔ'Øˆ
÷) 4Õ3ús   ÄHÈ
Hc                ó¶   • [         R                  R                  R                  U 5      nSU S[         R                  R                  R	                  SU5       S3$ )Nútl.load(ú + Úload_seed_offsetr¡  )rL   rÅ   rÂ  ÚinputÚseed_offset)rÈ   rº   rÚ   s      r†   rÛ  ÚTritonKernelOverrides.load_seed®  sI   € ä�h‰h�m‰m×!Ñ! $Ó'ˆà�s�e˜3œqŸx™xŸ}™}×8Ñ8Ð9KÈVÓTÐUÐUVÐWð	
rˆ   c                ó*  • SU  S3n[         R                  R                  R                  U5      =n(       a  U$ [         R                  R                  R	                  U R
                  U R                  S9n[         R                  R                  R	                  [        R                  U R                  S9n[         R                  R                  R                  U SU SU  S35        [         R                  R                  R                  XU45        X44$ )Nzfrexp(r¡  r0  r   z = triton_helpers.frexp()rL   rÅ   rÆ   Útry_getÚnewvarrƒ  rÉ   rœ   rŒ  r5  Ú	writelineÚput)rN  Ú	cache_keyÚcse_valÚmantissaÚexponents        r†   ÚfrexpÚTritonKernelOverrides.frexpµ  sÑ   € à˜Q˜C˜q�Mˆ	Ü—h‘h—l‘l×*Ñ*¨9Ó5Ð5ˆ7Õ5ØˆNä—8‘8—<‘<×&Ñ&¨Q¯W©W¸A¿G¹GÐ&ÐDˆÜ—8‘8—<‘<×&Ñ&¬U¯[©[ÀÇÁÐ&ÐHˆÜ	�‰×Ñ×"Ñ"Øˆj˜˜8˜*Ð$<¸Q¸C¸qÐAô	
ô 	
�‰�‰×Ñ˜¨xÐ$8Ô9ØÐ#Ð#rˆ   c                ó   • [         erš   ra  )rÈ   rƒ   rª  Ú
extra_metas       r†   Úpartial_accumulateÚ(TritonKernelOverrides.partial_accumulateÃ  s
   € ô "Ð!rˆ   r‚   )
rÈ   rÔ   rƒ   rÔ   rª  rT   rq  údict[str, Any]r¦   r§   )r›   r¨   r©   rª   r«   rš  r­   r6  Úcacher'  rî  rb  r$  rg  rÛ  rn  rr  r®   r¥  r¦  s   @r†   r3  r3    sÖ   ø† ñõ(ð Ø‡_�_ñ"4ó ó ð"4ðH ñ?ó ð?ð ñ;ó ð;ðz ñ,ó ð,ð\ ñ
ó ð
ð ñ$ó ð$ð ð"Øð"àð"ð ð"ð #ð	"ð
 
ó"ó ö"rˆ   r3  c                  óV   • \ rS rSr% SrS\S'   S\S'   SS jrSS	.SS
 jjrS rS r	Sr
g)ÚHelperFunctionsiÎ  z#An ordered set of helper functions.zdict[str, str]Ú_templates_seenú	list[str]Úfinalized_helpersc                ó    • 0 U l         / U l        g rš   ©rx  rz  r  s    r†   rš  ÚHelperFunctions.__init__Ô  s   € Ø!ˆÔØ!#ˆÕrˆ   Ú_triton_helper_fn©Ú	base_namec               óä   • U R                   R                  U5      nUb  U$ U [        U R                  5       3nX@R                   U'   U R                  R	                  UR                  US95        U$ )a  This accepts a function definition with the function name
left as a format specifier e.g.

    @triton.jit
    def {name}(arg0, arg1):
        return arg0 + arg1

We add the templated code to the function set and return the name
assigned to that function.

)rÈ   )rx  rè  r�   rz  rØ  rÃ  )r  Útemplate_coder€  Úexisting_namerÈ   s        r†   ÚaddÚHelperFunctions.addØ  su   € ð ×,Ñ,×0Ñ0°Ó?ˆØÑ$à Ð à�œS ×!7Ñ!7Ó8Ð9Ð:ˆØ.2×Ñ˜]Ñ+Ø×Ñ×%Ñ% m×&:Ñ&:ÀÐ&:Ð&EÔFØˆrˆ   c                ó,   • [        U R                  5      $ rš   )Úiterrz  r  s    r†   Ú__iter__ÚHelperFunctions.__iter__î  s   € Ü�D×*Ñ*Ó+Ð+rˆ   c                ó    • U R                   U   $ rš   )rz  )r  r…  s     r†   Ú__getitem__ÚHelperFunctions.__getitem__ñ  s   € Ø×%Ñ% cÑ*Ð*rˆ   r|  N©r¦   r§   )r‚  rÔ   r¦   rÔ   )r›   r¨   r©   rª   r«   r¬   rš  r„  rˆ  r‹  r®   r‚   rˆ   r†   rw  rw  Î  s+   ‡ Ù-à#Ó#Ø Ó ô$ð 4G÷ ò,,õ+rˆ   rw  c                  ó†  • \ rS rSr% Sr\R                  " \S9rS\	S'   \R                  " \S9r
S\	S'   \R                  " \S9rS\	S'   \R                  " \S9rS\	S'   \R                   " S	 S
5      5       r\R                   " S S\5      5       r\R                   " S S\5      5       rSS jr      SS jrSS jrSrg)r*  iõ  zE
Class representing ND block dimensions, for block pointer analysis.
)Údefault_factoryrO  rÉ   r;  r>  rB  c                  ó    • \ rS rSr% S\S'   S\S'   \R                  " SS9rS\S'   S r\	S	 5       r
\\      SS
 j5       5       rS rS rSrg)r4  é   r-  Úoriginal_stridesr¨  F)Úinitr©  c                óÌ  • [        U R                  5      S:”  d   e[        U R                  5      [        U R                  5      :X  d   e[        [	        [        U R                  5      5      5      nU R                  U:H  U l        [        U R                  5       VVs0 s H  u  p#X2_M	     nnn[	        [        U5      5       Vs/ s H  nXB   PM	     snU l        g s  snnf s  snf rx  )r�   r’  r¨  ra  rc  Ú_is_identityÚ	enumerater©  )r  Úidentity_sort_idxrU  r*  Úsorted_dims_by_strides_maps        r†   Ú__post_init__Ú*BlockParameters.StrideSorter.__post_init__  sÌ   € Ü�t×,Ñ,Ó-°Ó1Ð1Ð1Ü�t—}‘}Ó%¬¨T×-BÑ-BÓ)CÓCÐCÐCä $¤U¬3¨t×/DÑ/DÓ+EÓ%FÓ GÐØ $§¡Ð1BÑ BˆDÔô <EÀTÇ]Á]Ô;SÔ)TÒ;S±4°1¨!ª$Ñ;SÐ&Ñ)Tô œsÐ#=Ó>Ô?ó$â?�Að +Ô-Ù?ñ$ˆDÕ ùó *Uùò$s   ÂCÃC!c                ó   • U R                   $ rš   )r•  r  s    r†   r¦  Ú(BlockParameters.StrideSorter.is_identity  s   € à×$Ñ$Ð$rˆ   c                ó   • g)zBCreate a `StrideSorter` that can be used to sort block parameters.Nr‚   ©rŸ   r’  rM  s      r†   rl  Ú#BlockParameters.StrideSorter.create  s   � rˆ   c                óp   • U R                   (       d  U R                   Vs/ s H  o!U   PM	     sn$ U$ s  snf rš   )r¦  r¨  ©r  ÚattrrU  s      r†   ÚsortÚ!BlockParameters.StrideSorter.sort  s0   € Ø×#×#Ø)-¯ªÓ7ª A˜Qœ©Ñ7Ð7ØˆKùò 8ó    3c                óp   • U R                   (       d  U R                   Vs/ s H  o!U   PM	     sn$ U$ s  snf rš   )r¦  r©  r¡  s      r†   r£  Ú#BlockParameters.StrideSorter.revert$  s4   € Ø×#×#Ø)-×)=Ò)=Ó>Ò)= A˜QœÑ)=Ñ>Ð>ØˆKùò ?r¥  )r•  r©  N©r’  zlist[Union[int, sympy.Expr]]rM  rp   r¦   r4  )r›   r¨   r©   rª   r¬   ÚdataclassesÚfieldr©  r™  r'  r¦  r­   r   rl  r£  r£  r®   r‚   rˆ   r†   ÚStrideSorterÚBlockParameters.StrideSorter   s   ‡ à#Ó#ØÓØ%0×%6Ò%6¸EÑ%Bˆ˜ÓBò	ð 
ñ	%ó 
ð	%ð 
Ø	ð	UØ#?ð	UØLTð	Uà)ó	Uó 
ó 
ð	Uò
	õ
	rˆ   r«  c                  óH   ^ • \ rS rSrU 4S jr\      SS j5       rSrU =r$ )Ú$BlockParameters.IdentityStrideSorteri)  c                ó"   >• [         TU ]  5         g rš   ©r™  r™  ©r  rœ  s    €r†   r™  Ú2BlockParameters.IdentityStrideSorter.__post_init__+  ó   ø€ Ü‰GÑ!Õ#rˆ   c           
     óF   • U " U[        [        [        U5      5      5      S9$ )N©r’  r¨  )ra  rc  r�   rž  s      r†   rl  Ú+BlockParameters.IdentityStrideSorter.create.  s'   € ñ Ø!1Üœe¤CÐ(8Ó$9Ó:Ó;ñð rˆ   r‚   r¨  )	r›   r¨   r©   rª   r™  r­   rl  r®   r¥  r¦  s   @r†   ÚIdentityStrideSorterr®  )  s4   ø† õ	$ð 
ð	Ø#?ð	ØLTð	à)ó	ó 
ö	rˆ   r·  c                  óL   ^ • \ rS rSrSrU 4S jr\      SS j5       rSrU =r	$ )Ú+BlockParameters.TensorDecriptorStrideSorteri7  zD
Sorts BlockParameters dimensions with strides in descending order.
c                ó"   >• [         TU ]  5         g rš   r°  r±  s    €r†   r™  Ú9BlockParameters.TensorDecriptorStrideSorter.__post_init__=  r³  rˆ   c                óš   • [        [        [        U5      5      5      n [        R                  " X!SS9nU " UUS9$ ! [
         a    Un Nf = f)aV  
If the strides are not all known constants or if the strides are already
sorted in descending order, return identity sort.

For example if block_shape @ strides is [ZBLOCK, XBLOCK, YBLOCK] @ [8, 1, 16]
The indices to sort the strides in descending order will be [2, 0, 1].
The indices to revert back to the original order will be [1, 2, 0].
T)Úreverserµ  )ra  rc  r�   r%   r_  r`  )rŸ   r’  rM  Úidentity_sortr¨  s        r†   rl  Ú2BlockParameters.TensorDecriptorStrideSorter.create@  s^   € ô !¤¤sÐ+;Ó'<Ó!=Ó>ˆMð	)ô !×,Ò,¨YÐRVÑW�ñ
 Ø!1Ø!ñð øô	 "ó )à(’ð)ús   Ÿ; »A
Á	A
r‚   r¨  )
r›   r¨   r©   rª   r«   r™  r­   rl  r®   r¥  r¦  s   @r†   ÚTensorDecriptorStrideSorterr¹  7  s9   ø† ñ	õ	$ð 
ð	Ø#?ð	ØLTð	à)ó	ó 
ö	rˆ   rÀ  c                óŒ   • [        U 5      n[        S X4 5       5      u  p4U" S0 U Vs0 s H  oUX5   XE   -   _M     snD6$ s  snf )z 
Concatenates block parameters.
c              3  óN   #   • U  H  n[         R                  " U5      v •  M     g 7frš   )r©  Úasdictr:  s     r†   r  Ú*BlockParameters.__add__.<locals>.<genexpr>b  s   é € ÐB²M¨q”[×'Ò'¨×*Ð*²Mùs   ‚#%r‚   )rA  r×   )r  rf  rŸ   rD  rE  rÂ  s         r†   Ú__add__ÚBlockParameters.__add__]  sL   € ô �4‹jˆÜÑB°T±MÓBÓB‰ˆÙÑ9±aÓ8²a¨s˜1™6 A¡F™?Ò*±aÑ8Ñ9Ð9ùÒ8s   ©Ac                óê   • UR                  U R                  US9n[        S0 [        R                  " U 5      R                  5        VVs0 s H  u  pEXCR                  U5      _M     snnD6nXc4$ s  snnf )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.
)rM  r‚   )rl  r>  r*  r©  rÃ  r¼  r£  )r  rL  rM  r5  rÂ  rÃ  r+  s          r†   rT  Ú,BlockParameters.maybe_sort_with_stride_ordere  s}   € ð *×0Ñ0°·±ÈÐ0ÐSˆÜ ñ 
ô !,× 2Ò 2°4Ó 8× >Ñ >Ô @ôâ @‘H�Cð ×'Ñ'¨Ó,Ò,Ù @òñ
ˆð Ð$Ð$ùós   Á	A/c                ó¬   ^• U4S jn[        S0 [        R                  " U 5      R                  5        VVs0 s H  u  p4X2" U5      _M     snnD6$ s  snnf )z1
Remove dimensions where removable_dims is True.
c                ód   >• [        U T5       VVs/ s H  u  pU(       a  M  UPM     snn$ s  snnf rš   )r¤  )ÚitÚitemÚis_removableÚremovable_dimss      €r†   Úfilter_dimsÚ0BlockParameters.remove_dims.<locals>.filter_dims{  s7   ø€ ô +.¨b°.Ô*Aôâ*AÑ&�DÜ#÷ Ù*Aòð ùó s   ‘,¢,r‚   )r*  r©  rÃ  r¼  )r  rÎ  rÏ  rÂ  rÃ  s    `   r†   rS  ÚBlockParameters.remove_dimsv  sS   ø€ õ
	ô ñ 
Ü5@×5GÒ5GÈÓ5M×5SÑ5SÔ5UÔVÒ5U©¨ˆs�K Ó$Ò$Ñ5UÒVñ
ð 	
ùÛVs   ¶Ar‚   N)rf  r*  r¦   r*  )rL  ztype[StrideSorter]rM  rp   r¦   z4tuple[BlockParameters, BlockParameters.StrideSorter])rÎ  r1  r¦   r*  )r›   r¨   r©   rª   r«   r©  rª  ra  rÉ   r¬   r;  r>  rB  Ú	dataclassr«  r·  rÀ  rÅ  rT  rS  r®   r‚   rˆ   r†   r*  r*  õ  sé   ‡ ñð *×/Ò/ÀÑE€EÐÓEØ$/×$5Ò$5ÀdÑ$K€KÐ!ÓKØ +× 1Ò 1À$Ñ G€GÐÓGØ +× 1Ò 1À$Ñ G€GÐÓGà×Ñ÷&ð &ó ð&ðP ×Ñô˜|ó ó ðð ×Ñô# ló #ó ð#ôJ:ð%Ø!3ð%Ø@Hð%à	=ô%÷"
rˆ   r*  c                  ó4   • \ rS rSrSrS rS	S jrS rS rSr	g)
Ú"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                óš   • Xl         / U l        / U l        [        R                  " [        R
                  5      U l        SU l        SU l        g rx  )	rÂ  Úcurrent_loopÚ
prior_loopÚcollectionsÚdefaultdictÚdequeÚready_for_reuseÚ
loop_countÚstore_count)r  rÂ  s     r†   rš  Ú+CooperativeReductionWorkspaceCache.__init__�  s>   € ØŒ	ØˆÔØˆŒÜ*×6Ò6´{×7HÑ7HÓIˆÔØˆŒØˆÕrˆ   c                óä   • U R                   R                  U5      nU(       a  UR                  5       $ U R                  R	                  US5      u  p4nU R
                  R                  XU45        X54$ rŒ  )rÛ  rè  ÚpopleftrÂ  Ú	workspacerÖ  rØ  )r  ÚnbytesÚcachedÚws_namerW  Ú	ws_offsets         r†   ÚallocateÚ+CooperativeReductionWorkspaceCache.allocate•  se   € Ø×%Ñ%×)Ñ)¨&Ó1ˆÞØ—>‘>Ó#Ð#Ø $§	¡	× 3Ñ 3°F¸EÓ BÑˆ�IØ×Ñ× Ñ  &°9Ð!=Ô>ØÐ#Ð#rˆ   c                óÈ   • U R                    H%  u  pnU R                  U   R                  X#45        M'     U R                  U l         / U l        U =R                  S-  sl        g ©NrN   )r×  rÛ  rØ  rÖ  rÜ  )r  râ  rä  rå  s       r†   Úon_loop_endÚ.CooperativeReductionWorkspaceCache.on_loop_end�  sT   € à*.¯/¬/Ñ&ˆF˜YØ× Ñ  Ñ(×/Ñ/°Ð0DÖEñ +:à×+Ñ+ˆŒØˆÔØ�Š˜1ÑŽrˆ   c                óH   • U R                   nU =R                   S-  sl         U$ ré  )rÝ  )r  Úpriors     r†   Úincrement_store_countÚ8CooperativeReductionWorkspaceCache.increment_store_count¥  s#   € Ø× Ñ ˆØ×Ò˜AÑÕØˆrˆ   )rÂ  rÖ  rÜ  r×  rÛ  rÝ  N)râ  rí   )
r›   r¨   r©   rª   r«   rš  ræ  rê  rî  r®   r‚   rˆ   r†   rÔ  rÔ  ‡  s   † ñò
ô$òõrˆ   rÔ  c                  ó,   • \ rS rSr% S\S'   S rS rSrg)ÚFixedTritonConfigi«  zdict[str, int]r"   c                ó    • U R                   U   $ rš   ©r"   ©r  rÌ  s     r†   r‹  ÚFixedTritonConfig.__getitem__¯  s   € Ø�{‰{˜4Ñ Ð rˆ   c                ó   • XR                   ;   $ rš   ró  rô  s     r†   Ú__contains__ÚFixedTritonConfig.__contains__²  s   € Ø—{‘{Ñ"Ð"rˆ   r‚   N)r›   r¨   r©   rª   r¬   r‹  r÷  r®   r‚   rˆ   r†   rñ  rñ  «  s   ‡ àÓò!õ#rˆ   rñ  c                  ó"   • \ rS rSrSrSS jrSrg)Ú	TritonCSEi¶  zy
Subclasses CSE to apply the current load mask to the cache key to avoid CSEing
variables across separate masked blocks.
c                ób   • [         R                  R                  =n(       a  XR                  4$ U$ rš   )rL   rÅ   Ú
_load_maskrÈ   )r  rj  r  s      r†   Úaugment_keyÚTritonCSE.augment_key¼  s*   € Ü—8‘8×&Ñ&Ð&ˆ4Õ&ØŸy™yÐ)Ð)àÐrˆ   r‚   N)rj  rÔ   r¦   zUnion[str, tuple[str, str]])r›   r¨   r©   rª   r«   rý  r®   r‚   rˆ   r†   rú  rú  ¶  s   † ñ÷
rˆ   rú  c                  ór   • \ rS rSr% SrS\S'   S\S'   S\S'   S\S	'   S
 r  SS jr    SS jrSS jr	Sr
g)ÚTMACompatibilityCheckeriÃ  zG
Checks if the TMA API can be used for load / store triton operations.
ÚTritonKernelrÅ   r¤  rƒ  r¥   r­  Úforcec                ó   • SU l         g )Nz2Cannot use TMA descriptor for load / store since: ©Úfailed_debug_prefixr  s    r†   r™  Ú%TMACompatibilityChecker.__post_init__Î  s
   € Ø#WˆÕ rˆ   c                ó„  • U R                   (       a  g[        R                  R                  5       R                  S:X  a:  [
        R                  R                  5       S   S:¼  a  [        R                  (       d,  [        R                  R                  5       R                  S:X  a.  [        R                  R                  (       a  [        5       (       d!  [        R                  SU R                  5        gU R                   (       a<  U R"                  R$                  (       a!  [        R                  SU R                  5        gg)	NTrH  r   é	   Úxpuz}%s Requires triton>=3.4.0, a CUDA device with cc>=9.0 and `use_tensor_descriptor` and `assume_aligned_inputs` options enabledFz/%s stores with `no_x_dim` cannot load 16 bytes.)r  rL   rP  Úget_current_device_or_throwrA  rœ   rH  Úget_device_capabilityr"   Úassume_aligned_inputsrµ   Úuse_tensor_descriptorr   rø  Údebugr  r­  rÅ   rV  r  s    r†   Úcan_use_tmaÚ#TMACompatibilityChecker.can_use_tmaÒ  sØ   € ð �:�:Øô —G‘G×7Ñ7Ó9×>Ñ>À&ÓHÜŸ
™
×8Ñ8Ó:¸1Ñ=ÀÓBÜ×4×4ä—7‘7×6Ñ6Ó8×=Ñ=ÀÓFä—‘×3×3Ü)×+Ñ+ô �I‰Ið[ð ×(Ñ(ôð ð �>�>˜dŸk™k×2×2Ü�I‰IØAØ×(Ñ(ôð àrˆ   c           
     ó4  • U R                   (       aD  UR                   Vs/ s H,  n[        R                  R                  R                  U5      PM.     nnOUR                  n[        R                  R                  R                  US   [        R                  " S5      5      (       d"  [        R                  SU R                  U5        gU R                  R                  nUSS  H‹  n[        R                  R                  R                  [        XT-  S[        R                  " S5      5      [        R                  " S5      5      (       a  Mi  [        R                  SU R                  UU5          g   UR                  S   n[        R                  R                  R                  U[        R                  " S5      5      (       a,  [        R                  S	U R                  UR                  5        gSnSnUR                    H1  n	["        R$                   H  n
['        Xš5      (       d  M  U	nU
n  M/     M3     U(       a  U(       d   U S
["        R$                   35       eU R(                  R*                  (       GaR  U R,                  (       Gd@  U["        R.                  ;   Ga+  [0        U   nSnU R(                  R2                   H4  nUR4                  (       d  M  UR6                  U:X  d  M(  UR8                  n  O   Uc-  [        R                  SU R                  UUR                  5        gU R(                  R;                  U5      nUR=                  X~05      U-  n[        R                  R                  R?                  U[        R                  " S5      5      (       d-  [        R                  SU R                  UR                  U5        g g  S       SS jjnXd-  S-
  nURA                  [B        U5      RA                  [        U5      n[E        [G        [        RH                  " UUS5      5      5      nUU R(                  RK                  [0        U   5      :”  a#  [        R                  SU R                  UU5        gU R(                  RM                  U5      n	U R(                  RN                  (       aZ  UU R(                  RN                  U	   :”  a;  [        R                  SU R                  U	U R(                  RN                  U	   U5        g g[Q        UU R(                  RR                  RU                  U	S5      5      U R(                  RR                  U	'    gs  snf ! [V         a.    [        R                  SU R                  UR                  5         gf = f)z‘
Check if the block parameters are valid for TMA.
If force, we allow relying on symbolic hints equivalent
to what we check for Triton templates.
rK  rN   z=%s TMA API requires innermost stride to be 1. Strides are: %sFNé   r   zU%s TMA API requires outer strides to be 16 byte aligned. Dtype bytes: %d, strides: %sz>%s innermost block shape cannot load 16 bytes. Block shape: %sz, expr must contain a single block type from zN%s could not find reduction range tree for innermost prefix %s Block shape: %szj%s persistent reduction innermost block shape cannot load 16 bytes. Block shape: %s, persistent RBLOCK: %dc                ó$   • X-  nU(       a  X2-  nU$ rš   r‚   )rN  rù  r  r  s       r†   Úindexing_div_repÚQTMACompatibilityChecker.are_block_parameters_compatible.<locals>.indexing_div_repv	  s   € ð
 ™%�CÞØ!™g˜Ø�Jrˆ   zC%s the minimum block size to satisfy expression %s is too large: %dzT%s For block %s, fixed config block size %d is smaller than the minimum required: %dz?%s innermost block shape cannot load 16 bytes. Block params: %sTrš   )rN  rí   rù  rí   r  zOptional[sympy.Expr]r¦   rí   ),r  r>  rL   rP  rH  Úsymbolic_hintrQ  r‹   r�   rø  r  r  rƒ  r‹  r   r;  r�   r¸   rÐ   r   rÅ   Úpersistent_reductionr­  rð   r   re  r\  rÓ   ÚnumelÚ_get_persistent_RBLOCKÚsubsÚstatically_known_geqÚreplacer   r5   rî  ÚnsolveÚ	max_blockr§  Úfixed_configr{   Útma_min_block_sizesrè  r{  )r  Úblock_paramsÚstr>  Úelement_sizeri  Úinnermost_block_shapeÚinnermost_block_typeÚinnermost_block_symtÚblock_type_strÚ
block_symtÚinnermost_tree_prefixÚ
tree_numelr‚  Úpersistent_rblockÚinnermost_block_bytesr  Ú
solve_exprÚsolve_expr_simplifiedÚmin_block_sizes                       r†   Úare_block_parameters_compatibleÚ7TMACompatibilityChecker.are_block_parameters_compatibleú  s  € ð �:�:à=I×=QÒ=QóÚ=Q°r”—‘× Ñ ×.Ñ.¨rÖ2Ñ=Qð ð ˆGð #×*Ñ*ˆGô �w‰w×Ñ×7Ñ7¸À¹ÄUÇ]Â]ÐSTÓEU×VÑVÜ�I‰IØOØ×(Ñ(Øôð
 à—z‘z×*Ñ*ˆØ˜c˜r“lˆFÜ—7‘7×#Ñ#×;Ñ;Ü Ñ 5°q¼%¿-º-ÈÓ:KÓLÜ—’˜aÓ ÷ó ô —	‘	ØkØ×,Ñ,Ø Øô	ñ ñ #ð" !-× 8Ñ 8¸Ñ <Ðô �7‰7×Ñ×3Ñ3Ø!¤5§=¢=°Ó#3÷
ñ 
ô �I‰IØPØ×(Ñ(Ø×(Ñ(ôð
 à#ÐØ#ÐØ3×@Ô@ˆNÜ+×7Ô7�
Ü! .×=Ó=Ø+9Ð(Ø+5Ð(Úó	 8ñ Aö $Ö(<ð 	
Ø$Ð%Ð%QÔR_×RkÑRkÐQlÐmó	
Ð<ð �K‰K×,×,Ð,Ø—N—N�NØ$¬×(EÑ(EÔEô %/Ð/CÑ$DÐ!ØˆJØ—[‘[×,Ô,�Ø—>—>‘> a§h¡hÐ2GÕ&GØ!"§¡�JÙñ -ð Ñ!ô
 —	‘	ØdØ×,Ñ,Ø)Ø ×,Ñ,ô	ð Ø $§¡× BÑ BÀ:Ó NÐà%×*Ñ*Ð,@Ð+TÓUØñð "ô —7‘7×#Ñ#×8Ñ8Ø%¤u§}¢}°RÓ'8÷ñ ô —	‘	ð AØ×,Ñ,Ø ×,Ñ,Ø%ô	ð ððf ðIBð
 /3ðØ!ðà!ðð ,ðð  õ	ð 3ÑAÀBÑF�
à(2×(:Ñ(:ÜÐ.ó)ç‘'œ/Ð+;Ó<ð &ô "1ÜÜŸšØ1Ø0Øóóó"�ð " D§K¡K×$9Ñ$9ÜÐ3Ñ4ó%ó ô —I‘IØ]Ø×0Ñ0Ø-Ø&ô	ð !à!%§¡×!9Ñ!9Ð:NÓ!O�à—;‘;×+×+Ø%¨¯©×(@Ñ(@ÀÑ(PÓPÜŸ	™	ð<à ×4Ñ4Ø*Ø ŸK™K×4Ñ4°^ÑDØ*ôð  %ð Qð4 ô GJØ&ØŸ™×7Ñ7×;Ñ;¸NÈAÓNóG�D—K‘K×3Ñ3°NÒCð ùòiøôX ó Ü—	‘	ØUØ×,Ñ,Ø ×,Ñ,ôñ
 ðús&    3UÏB2U ÒBU ÔAU Õ5VÖVc                ó   • U R                   $ )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.
)r  r  s    r†   r7  Ú TMACompatibilityChecker.can_liftº	  s   € ð �z‰zÐrˆ   r  Nr%  )r!  r*  r¦   r¥   )r›   r¨   r©   rª   r«   r¬   r™  r  r0  r7  r®   r‚   rˆ   r†   r   r   Ã  sS   ‡ ñð ÓØÓØƒOØƒKòXð&à	ô&ðP~à%ð~ð 
ô~÷@
rˆ   r   c                  óä  ^ • \ rS rSr% Sr\rS\S'   \r	S\S'   Sr
\r\rS\S	'   \rS
\S'   SrS\S'        Sb         ScU 4S jjjr\SdS j5       r\SdS j5       rSeS jrSdS jrS rS rS rS rSdS jrS r\SfS j5       rSSSSSS.     SgS jjr  Sh       SiS jjr!ShS jr"        SjS  jr#S! r$S"r%S#r&\S$ 5       r'SS%.S& jr(S' r)SkS( jr*    SlS) jr+SmS* jr, Sn         SoS+ jjr-SpS, jr.S- r/SqS. jr0  Sr               SsS/ jjr1SfS0 jr2StS1 jr3      SuS2 jr4          SvS3 jr5  SwS4 jr6SwS5 jr7S6 r8S7 r9S8 r:S9 r;  SxS: jr<      SyS; jr=      SzS< jr>        S{S= jr?          S|S> jr@SkS? jrAS@ rBS}SA jrCS~SB jrDSC rESD rF\G\H" S5      SfSE j5       5       rI\GSF 5       rJ\GSG 5       rKSnSfSH jjrL\SI 5       rM\SJ 5       rNSK rOSSL jrPSM rQ S€     S�SN jjrRSpSO jrSS‚SP jrTSƒSQ jrUS„SR jrV      S…SS jrWS„ST jrXS†SU jrYS‡SV jrZSˆSW jr[SdSX jr\S‰SY jr]\^S}SZ j5       r_SŠS[ jr`S‹S\ jra\^SŒS] j5       rbS�S^ jrcSŠS_ jrd      SŽS` jreSarfU =rg$ )�r  iÇ	  z\A class to represent a triton kernel and helpers to generate
triton kernel programmatically
rw  Úhelper_functionszCallable[[sympy.Expr], str]ÚkexprTztype[BlockPtrOptions]Úblock_ptr_options_clsztype[TensorDescriptorOptions]Útensor_descriptor_options_clsNzOptional[bool]Ú3transpose_discontiguous_tensor_descriptors_overrideFc                óì  >• X0l         X@l        X`l        [        TU ]  " U40 UD6  [        U R                  U R                  5      U l        0 U l	        [        5       U l        [        5       U l        [        5       U l        [        [           " 5       U l        X l        [$        R&                  " 5       U l        [*        [,        [,        4   " 5       U l        [1        5       U l        [4        R6                  " [*        5      U l        [*        [,        [:        4   " 5       U l        XPl        [4        R@                  " 5       U l!        SU l"        SU l#        [        [H           " 5       U l%        S U l&        U RN                  (       a  U RQ                  U RR                  5        U RT                  (       a  U RW                  5         U RY                  5         U RT                  (       a  U R[                  5         SU l.        / U l/        g )Nr   F)0Úoptimize_maskr  Úis_combo_kernelr™  rš  rú  Únewvar_prefixÚsuffixrÆ   Úprologue_cacherV   ÚprologueÚpost_loop_combineÚpost_loop_storer   r	   Úoutside_loop_varsrÚ  r½  ÚcountÚblock_ptr_idÚdictrÔ   Úblock_ptr_to_bufferrw  r5  rØ  rÙ  Úpointer_advancementsrî  r   Úhint_overrideÚCounterÚ_load_countsÚ_pdl_load_indexÚ_pdl_has_waitr/   Úautotune_hintsÚtriton_metarY  Úcodegen_reduction_numelsre  Úcooperative_reductionÚinit_cooperative_reductionÚcodegen_range_treeÚinit_cooperative_reduction_maskÚhas_load_with_contiguous_rdimÚstores_with_contiguous_rdim)	r  ÚtilingrÚ  r;  r  rI  r<  r   rœ  s	           €r†   rš  ÚTritonKernel.__init__×	  sx  ø€ ð $1ÔØ(ÔØ%4ÔÜ‰Ò˜Ñ* 6Ò*Ü˜T×/Ñ/°·±Ó=ˆŒà.0ˆÔÜ(6Ó(8ˆŒÜ1?Ó1AˆÔÜ/=Ó/?ˆÔÜ!+¬C¢Ó!2ˆÔØ#6Ô Ü%ŸOšOÓ-ˆÔÜ#'¬¬S¨¢>Ó#3ˆÔ Ü /Ó 1ˆÔä×#Ò#¤DÓ)ð 	Ô!ô $(¬¬S¨¢>Ó#3ˆÔ Ø*ÔÜ6A×6IÒ6IÓ6KˆÔØ ˆÔØ"ˆÔô )¬Ò6Ó8ˆÔØ59ˆÔà× × Ø×)Ñ)¨$¯)©)Ô4à×%×%Ø×+Ñ+Ô-à×ÑÔ!à×%×%Ø×0Ñ0Ô2à-2ˆÔ*à68ˆÕ(rˆ   c                ó®  • [         R                  (       d   [         R                  R                  (       d  gU R                  nU Vs/ s H&  n[        U[        R                  5      (       d  M$  UPM(     nn[        U5      S:X  a  g [        R                  R                  R                  XU5      n[        S U 5       5      $ s  snf ! [         a     gf = f)NFr   c              3  ó*   #   • U  H	  oS :H  v •  M     g7fr›  r‚   )rö   ri  s     r†   r  Ú4TritonKernel._has_stride1_on_rdim.<locals>.<genexpr>&
  s   é € Ð9ª[ 6˜Q–;ª[ùó   ‚)r"   ÚdeterministicrG  Úforce_filter_reduction_configsr�   r   r¸   rð   r�   rL   rP  rH  Ústride_varsÚZeroDivisionErrorrz   )rþ   Úsupport_varsrÚ   Úreduce_varsr_  s        r†   Ú_has_stride1_on_rdimÚ!TritonKernel._has_stride1_on_rdim
  sº   € ô × × ¤F×$7Ñ$7×$V×$VàØ×)Ñ)ˆñ $ó
â#�Ü˜c¤=×#@Ñ#@×A÷ Ù#ð 	ð 
ô ˆ{Ó˜qÓ Øð	ÜŸ'™'×*Ñ*×6Ñ6°uÈ<ÓXˆKô Ñ9©[Ó9Ó9Ð9ùò!
øô !ó 	Ùð	ús   Á#CÁ-CÂ*C Ã
CÃCc                óD   • [        S U R                   5       5      (       + $ )Nc              3  ó8   #   • U  H  n[        U5      v •  M     g 7frš   )rX   )rö   rÈ   s     r†   r  Ú>TritonKernel.has_store_with_contiguous_rdim.<locals>.<genexpr>*
  s   é € ð 
Ú0P¨Ô˜d×#Ð#Ò0Pùó   ‚)rR  rV  r  s    r†   Úhas_store_with_contiguous_rdimÚ+TritonKernel.has_store_with_contiguous_rdim(
  s'   € äñ 
Ø04×0PÒ0Pó
ó 
ô 
ð 	
rˆ   c                ó   • [        U5      $ rš   )rF   )r  rƒ  s     r†   Údtype_to_strÚTritonKernel.dtype_to_str.
  s   € Ü˜5Ó!Ð!rˆ   c                óz   • U R                   =(       a)    [        R                  R                  U R                  5      $ rš   )rY  rL   ÚchoicesÚ should_use_cooperative_reductionr[  r  s    r†   rp  Ú-TritonKernel.should_use_cooperative_reduction1
  s-   € Ø×$Ñ$÷ 
¬¯©×)SÑ)SØ�M‰Mó*
ð 	
rˆ   c                ó6  ^ • T R                   (       d   eT R                   H'  nUR                  c  M  U=R                  S-  sl        M)     T R                  S   nT R                  (       a  [        UT R                  S   5      nT R                  R                  U5      T l        [        T R                  5      T l
        T R                  R                  S5        [        U 4S jT R                   5       5      (       a  T R                  R                  S5        gg)z/One time setup code for cooperative reductions.NrN   rN  rñ   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)
            c              3  óv   >#   • U  H.  nUR                   (       d  M  TR                  U5      (       + v •  M0     g 7frš   )r\  r'  ©rö   rá   r  s     €r†   r  Ú:TritonKernel.init_cooperative_reduction.<locals>.<genexpr>R
  s4   øé € ð 
â(�Ø× Õ ó .�×'Ñ'¨Ó-×-Ñ-Ú(ùs   ƒ9œ9z>rsplit_end = tl.where(rsplit_end < rnumel, rsplit_end, rnumel))rQ  re  Úgrid_dimrZ  r  r   rÂ  Ú
semaphoresÚsemaphores_namerÔ  Ú%cooperative_reduction_workspace_cachere  Úsplicerz   rh  )r  rá   Ú	sem_counts   `  r†   rR  Ú'TritonKernel.init_cooperative_reduction6
  sî   ø€ à×)×)Ð)Ð)ð ×$Ô$ˆDØ�}‰}Ó(Ø—’ Ñ"—ñ %ð —K‘K Ñ$ˆ	Ø××Ü 	¨4×+<Ñ+<¸XÑ+FÓGˆIØ#Ÿy™y×3Ñ3°IÓ>ˆÔÜ5WØ�I‰Ió6
ˆÔ2ð 	�	‰	×Ñð	ô	
ô ô 
à×(Ò(ó
÷ 
ñ 
ð
 �I‰I×ÑØPõð
rˆ   c                ó.  • SnU R                   (       d  U S3nU R                  R                  SU 35        U R                  5       (       a  U R                  R	                  S5        g U R                   (       a   eU R                  R                  S5        g )Nz$tl.arange(0, RSPLIT_NEXT_POWER_OF_2)z	[None, :]zrsplit_arange = z½                if RSPLIT_IS_POWER_OF_2:
                    rsplit_mask: tl.constexpr = None
                else:
                    rsplit_mask = rsplit_arange < RSPLIT
                zSrsplit_mask = xmask if RSPLIT_IS_POWER_OF_2 else ((rsplit_arange < RSPLIT) & xmask))rV  re  rh  Ú_has_constant_xmaskrz  )r  Úrsplit_aranges     r†   rT  Ú,TritonKernel.init_cooperative_reduction_mask[
  s~   € Ø>ˆØ�}�}Ø,˜o¨YÐ7ˆMØ�	‰	×ÑÐ.¨}¨oÐ>Ô?à×#Ñ#×%Ñ%Ø�I‰I×Ñðõð —}—}Ð$Ð$Ø�I‰I×ÑØeõrˆ   c                óZ  • U R                    H}  nUR                  (       d  U R                  XR                  5        M1  U R                  (       d  MD  U R                  R                  UR                   SU R                  U5       35        M     U R                  (       a�  [        S U R                    5       5      (       aP  U R                  SSSS9nU R                  U5      nU R                  R                  SU R                  U5       35        g U R                  U R                  5        g g )Nzbase = c              3  ó8   #   • U  H  oR                   v •  M     g 7frš   )Úis_loop©rö   rá   s     r†   r  Ú2TritonKernel.codegen_range_tree.<locals>.<genexpr>}
  s   é € Ð=Ò,< D—<–<Ò,<ùrh  r$  Tr»   zrbase = )re  rƒ  Úiteration_ranges_codegen_headerre  rY  rh  rÓ   Úiteration_ranges_ranges_coderz   Ú_get_reduction_symbolsÚ_flatten_reduction_indicesrz  r§  Úcodegen_reduction_indices)r  rá   Úrn_basesÚrbases       r†   rS  ÚTritonKernel.codegen_range_treep
  s÷   € Ø×$Ô$ˆDà—<—<Ø×4Ñ4°T¿9¹9ÖEØ×&×&Ñ&ð —	‘	×#Ñ#Ø—{‘{�m 7¨4×+LÑ+LÈTÓ+RÐ*SÐTöñ %ð × × ÜÑ=¨D×,<Ò,<Ó=×=Ñ=à×6Ñ6Ø D°dð 7ð �ð ×7Ñ7¸ÓA�Ø—	‘	× Ñ  8¨D×,=Ñ,=¸eÓ,DÐ+EÐ!FÕGð ×.Ñ.¨t¯y©yÕ9ð !rˆ   c                ó   • g)z»
Indicate whether we need provide numel as arguments for the generated
kernel calls in the benchmark.

Should be true for pointwise/reduction kernels but false for triton
matmul kernels.
Tr‚   r  s    r†   Úneed_numel_argsÚTritonKernel.need_numel_argsˆ
  s   € ð rˆ   c                ó�   • U R                   =(       a4    [        R                  R                  U R                  U R
                  5      $ rš   )rY  rL   ro  Úshould_use_persistent_reductionr[  rQ  r  s    r†   r’  Ú,TritonKernel.should_use_persistent_reduction’
  s5   € Ø×$Ñ$÷ 
¬¯©×)RÑ)RØ�M‰M˜4×5Ñ5ó*
ð 	
rˆ   c                óÊ   • U R                   =(       aQ    [        U R                  5      U R                  S-   :H  =(       a%    U R                  =(       a    U R                  S   S:H  $ )NrN   rñ   )r  r�   rZ  rX  r  r  s    r†   Úwant_no_x_dimÚTritonKernel.want_no_x_dim—
  sY   € à×%Ñ%÷ 1Ü�D—K‘KÓ  D×$;Ñ$;¸aÑ$?Ñ?÷1à×!Ñ!÷1ð ×!Ñ! (Ñ+¨qÑ0ð		
rˆ   c                ó   • g)Nztl.device_assertr‚   r  s    r†   Úassert_functionÚTritonKernel.assert_functionŸ
  s   € à!rˆ   )Ú
copy_shapeÚdense_indexingÚoverride_maskrA  rB  c          
     óì  ^ ^^^^!^"^#^$• T R                  T5      mTR                  nSn[        5       m![        U[        R
                  " S5      S9 GH¡  n	[        U	[        R                  5      (       d   eU=(       d    [        U	[        R                  5      nU(       a  MQ  [        U	[        R                  5      (       a@  T R                  R                  U	R                      n
T!R#                  U
R$                  5        M°  [        U	[        R&                  [        R(                  [        R*                  [        R,                  [        R.                  [        R0                  45      (       a  GM  [        R2                   Vs/ s H  n[        X›5      (       d  M  [4        U   PM      nn[7        U5      S:X  a   [7        U5      S:X  d   SU	R                    35       eT!R9                  US    S35        GM¤     [:        R<                  R>                  =(       d    U=(       d    T R@                  SL=(       a    TS:g  nS	nSn[        5       nT RC                  5        HF  nURE                  URF                  5      (       a  S	nOSnUR9                  URH                   S35        MH     U(       a0  T RJ                  (       a  [:        R<                  RL                  (       d  T(       a­  TRO                  5       (       a˜  U(       d‘  T R@                  (       d€  [7        T!U-
  5      S:X  an  T RQ                  T5      (       dX  U(       aQ  T RR                  S
:X  aA        SS jm"      SU 4S jjm$      SU"U$4S jjm#SUUU!U#U U4S jjnU" 5       nUb  U$ SnSnT RU                  T5      nUU 4S jn[W        T5      (       Ga
  T(       d  [7        T RY                  5       5      S:X  a  U" 5       u  nnOL[[        S/[7        T RY                  5       5      -  5      n[]        S/[7        T RY                  5       5      -  5      nSU SU S3nT R^                  (       d  T R`                  (       a   [        U 4S jT Rb                   5       5      m!O
[        5       m!T R@                  (       a  T!R9                  T R@                  5        [e        UT!UUTUS9$ U(       Ga  U(       Gdþ  T Rf                  (       GaÖ  T Rh                  (       GaÄ  T!Rk                  5       nT R@                  (       a  UR9                  T R@                  5        [        / SQ5      nURm                  U5      (       d|  URo                  U5      nURq                  5       n[        U[r        5      (       d   eURu                  U5        UR#                  UR$                  5        URm                  U5      (       d  M|  S/[7        T RY                  5       5      -  nU HŒ  n[        U[Z        5      (       d   eT RC                  5        H^  nURw                  URH                  5      (       d  M%  URx                  n[        U[z        5      (       d   eT RY                  5       U   UU'   M`     MŽ     SSR}                  [        [Z        U5      5      -   S-   n[]        U5      nSU SU S3nO9U" 5       u  nnSU SU S3nUm!O#U(       d  T(       a  U" 5       u  nnSU SU S3nUm!Uc  U(       d  U(       a  U" 5       u  n nOSnU(       a  [        U/5      m!T R@                  (       a  T!R9                  T R@                  5        T R�                  T!5        [e        UT!UUTUS9$ s  snf ) z?
Compute the index and mask to pass to tl.load() or tl.store()
FrÈ   ©rÂ  r   rN   rÂ   r  NTútl.int32c                óÒ   • [         R                  " XR                  5       5      nUc  g[        UR                  /[
        R                  U5      /U/[
        R                  U5      /S9$ )zg
Matches expressions of the form:
    idx = s * xindex

This implies stride (s,), and shape (XBLOCK,).
N©rÉ   r;  r>  rB  )rO   Úmatch_affine_block_exprÚsymbolr*  r  r¸   rÕ   rë   )rþ   Ú
range_treeri  s      r†   Úmatch_affine_blockÚ1TritonKernel.indexing.<locals>.match_affine_blockú
  sj   € ô -×DÒDØ×,Ñ,Ó.ó�ð ‘>Øä&Ø%×+Ñ+Ð,Ü!.×!=Ñ!=¸jÓ!IÐ JØ#˜HÜ*×;Ñ;¸JÓGÐHñ	ð rˆ   c                óè  >^^• UR                  5       n[        R                  " S[        R                  " [        R
                  U/S9S9u  p4[        S[        UR                  5      U R                  [        X#5      5      U R                  [        X#U5      5      -   5      n[        R                  R                  mTR                  U 5      n TR                  UR                   5      n["        R$                  " XXe5      nUc  gUu  nn	n
["        R&                  " U5      n[        R                  R                  mTR)                  UR*                  5      m[-        UU4S jU 5       5      (       a  g[.        R1                  U5      n[3        UTR5                  US   5      5      /[7        USS USS 5       VVs/ s HD  u  pm[        R8                  " [3        UTR5                  U5      5      TR5                  U5      5      PMF     snn-   nU
 Vs/ s H#  n[;        Xò[.        R=                  U5      05      PM%     nn[?        U Vs/ s H  nTR5                  U5      PM     snUU	US	9$ s  snnf s  snf s  snf )
a  
Matches higher-dimensional blocks coming from FloorDiv and ModularIndexing.

Example expression to match:
   sN * ((rindex//(d1 * ... * d(N-1))))
       + s1 * ModularIndexing(rindex, 1, d1)
       + ...
       + s(N-1) * ModularIndexing(rindex, d1 * ... * d(N-2), d(N-1))

This iterates over a block of shape (dN, ..., d1) and stride
(sN, ..., s1). (d1,...,d(N-1)) and (s1,...,sN) are
wildcards that we match.

Note that dN does not appear in the expression, but we solve for it
using range tree numels and the other dims.
zdenom modulo)Úexclude©rŸ   r!   Nc              3  óŒ   >#   • U  H9  nTR                  UT5      (       + =(       a    TR                  U5      (       + v •  M;     g 7frš   )r�  Ústatically_known_power_of_2)rö   r  r  rH  s     €€r†   r  ÚETritonKernel.indexing.<locals>.match_mod_div_block.<locals>.<genexpr>]  sG   øé € ð ò ".˜ð !×=Ñ=¸eÀYÓOÔO÷ HØ$×@Ñ@ÀÓGÔGôHâ!-ùs   ƒAAr   rN   r¡  ) r£  r‹   Úsymbolsr6  r7  ÚWildr{   r�   rY  rD  r   r   rL   rP  rH  Úremove_precomputed_replacementsr  rO   Úmatch_mod_div_block_exprÚget_slice_numelsr  rÓ   rz   r¸   rÕ   r   rF  r¤  ÚMinrE   rë   r*  )rþ   r¤  rK  ÚdenomÚmoduloÚnum_dimsr  Úmatch_resultÚdimsr>  Úblock_index_exprsÚslice_numelsÚlinear_block_sizerg  r;  r’   rê   Údr  rH  r  s                     @@€r†   Úmatch_mod_div_blockÚ2TritonKernel.indexing.<locals>.match_mod_div_block  sq  ú€ ð& '×-Ñ-Ó/�	ô !&§¢Ø"Ü!×)Ò)¬%¯*©*¸y¸kÑJñ!‘�ô
 Øô ˜
×(Ñ(Ó)àŸ™¤H¨YÓ$>Ó?ØŸ+™+¤o°iÈÓ&OÓPñQó	�ô& Ÿ7™7×+Ñ+�Ø ×@Ñ@ÀÓG�Ø ×@Ñ@À×AQÑAQÓR�ä2×KÒKØ eó �ð  Ñ'Øð !ñ	ØØØ%ä2×CÒCÀDÓI�ô Ÿ7™7×+Ñ+�Ø ŸN™N¨:×+<Ñ+<Ó=�	Üõ ñ ".ó÷ ñ ð
  ô %2×$@Ñ$@ÀÓ$LÐ!äØ)Ø ×8Ñ8¸Àa¹ÓIóð1ô '*¨,°q°rÐ*:¸DÀÀ¸HÔ&Eôò 'F™
˜ô —I’IÜØ-¨x×/OÑ/OÐPUÓ/Vóð !×8Ñ8¸Ó=ö	ñ 'Fòñ1�ñ( !2ó	3ò !2˜ô Ø¬-×*HÑ*HÈÓ*TÐUöñ !2ð	 ð 3ô 'ÙHLÓMÊÀ1˜8×;Ñ;¸AÖ>ÉÑMØ +Ø#Ø)ñ	ð ùó%ùò3ùò Ns   Æ7AI$È*I*ÉI/c                ó‚   >• TT4 H6  n[         R                  " XR                  5       5      nU" X15      nUc  M4  Us  $    g)zE
Match a block indexing subexpression involving a single range tree.
N)rO   Úfactor_index_exprr£  )r’   r¤  Ú
match_funcÚfactored_index_exprÚmatchr¥  r¼  s        €€r†   Úmatch_block_subexprÚ2TritonKernel.indexing.<locals>.match_block_subexprŠ  sS   ø€ ð 'Ø'ó#�Jô +>×*OÒ*OØ×/Ñ/Ó1ó+Ð'ñ 'Ð':ÓG�EØÓ(Ø$šñ#ð rˆ   c            
     óî  >• [        TTR                  R                  5        V Vs0 s H  u  pXR                  _M     snn 5      nTR	                  5       nU Vs/ s H'  n[
        R                  " X$R                  5       5      PM)     nn[        S U 5       5      n[        5       n[        X55       H@  u  pH[        UR                  UR                  5      5      S:”  a    g T" X„5      n	U	c    g Xy-  nMB     U[        U5      -
  n
TR                  T5        [         R"                  R$                  (       a  TR&                  OTR(                  n[         R"                  R$                  (       a  Sn[        R*                  nO”[-        [.        T5      mTR1                  5       nTR2                   b  TR2                  nO[         R"                  R4                  n[7        TS5      (       a  UTS L-  nU(       a  [        R8                  O[        R*                  nUR;                  UU
UTTR<                  UUS9nU[>        :X  a1  [-        [.        T5      mTRA                  URB                  5      (       d  g U$ s  snn f s  snf )Nc              3  ó@   #   • U  H  oR                  5       v •  M     g 7frš   )r£  r„  s     r†   r  ÚBTritonKernel.indexing.<locals>.match_block_expr.<locals>.<genexpr>¯  s   é € Ð*QÂ[¸T¯;©;¯=¨=Â[ùó   ‚rN   FÚtemplate_out_shape)r+  r,  re  rû   rf  r7  rL  )"rE   Úrange_tree_nodesr¼  r’   rÒ   rO   Úget_subexpr_involving_symbolr£  r   r*  r¤  r�   Úintersectionr�   ÚsumÚfilter_masksr"   rµ   Úuse_block_ptrr7  r8  r·  r
   r   r7  r9  Ú)transpose_discontiguous_tensor_descriptorr³   rÀ  rl  r  r»  r0  r+  )r+  r‚  Úindex_relative_to_xyr_indexre  rá   Úindex_subexprsÚrange_symbolsr!  Úsubexprr+  rº   Úoptions_classr7  rL  Útranspose_contiguousÚoptionsrš  rþ   rû   rÃ  r  rB  s                   €€€€€€r†   Úmatch_block_exprÚ/TritonKernel.indexing.<locals>.match_block_exprž  sm  ø€ Ü.8Ø°$×2GÑ2G×2MÑ2MÔ2OÔPÒ2O©$¨!˜AŸv™všIÑ2OÒPó/Ð+ð #×5Ñ5Ó7�ñ !,ó	"ò !,˜ô (×DÒDØ3·[±[³]öñ !,ð	 ð "ô !+Ñ*QÁ[Ó*QÓ Q�Ü.Ó0�Ü%(¨Ö%E‘M�Dô ˜=×5Ñ5°g×6JÑ6JÓKÓLÈqÓPÙ#ñ 1°Ó?�FØ‘~Ù#Ø Ñ*’Lñ &Fð 5´s¸>Ó7JÑJ�ð ×!Ñ! )Ô,ô —}‘}×2×2ð ×.Ò.à×;Ñ;ð ô —=‘=×.×.Ø$�HÜ(7×(LÑ(LÑ%ä04Ü/Ð1Jó1Ð-ð  9×AÑAÓC�Hð ×PÑPØ#ñ$ð !×TÑTñ -ô
 #ŸM™M×SÑSð -ô ˜tÐ%9×:Ñ:Ø,°
À$Ð0FÑFÐ,ö 0ô (×CÒCä,×AÑAð &ð (×.Ñ.Ø'Ø$*Ø +Ø'Ø"&§.¡.Ø%Ø&7ð /ð �ð !Ô$;Ó;Ü04Ü/Ð1Jó1Ð-ð 5×TÑTØŸ™÷ñ ð  $à�ùó{ Qùò"s   ¥I,Á.I2c                 óÜ   >• T (       a<  [        T [        5      (       a  T  S3S 4$ SSR                  S T  5       5      -   S-   T 4$ TR                  5       [	        TR                  5       5      4$ )Nz.shaperÖ  r   c              3  ó8   #   • U  H  n[        U5      v •  M     g 7frš   ©rÔ   )rö   r3  s     r†   r  ÚATritonKernel.indexing.<locals>._get_expand_str.<locals>.<genexpr>  s   é € Ð*Fº:°a¬3¨q¯6¨6º:ùrh  r×  )rŠ   rÔ   r   Údense_size_strr×   rF  )rš  r  s   €€r†   Ú_get_expand_strÚ.TritonKernel.indexing.<locals>._get_expand_str  sk   ø€ ÞÜ˜j¬#×.Ñ.Ø(˜\¨Ð0°$Ð6Ð6à §¡Ñ*F¹:Ó*FÓ!FÑFÈÑLÈjÐXÐXà×*Ñ*Ó,¬e°D×4HÑ4HÓ4JÓ.KÐKÐKrˆ   rå  r   z, tl.int32)c              3  ó˜   >#   • U  H?  nUR                   (       a  M  TR                  U5      (       a  M.  UR                   S 3v •  MA     g7f©r  N)r\  r'  rÓ   rt  s     €r†   r  Ú(TritonKernel.indexing.<locals>.<genexpr>$  sA   øé € ð 'â 0˜Ø×,Õ,ó )à59×5LÑ5LÈT×5Ró )�t—{‘{�m 4Õ(Ú 0ùs   ƒA
œA
´A
)rÿ   )ÚxmaskÚymaskÚzmaskÚr0_maskrÃ   rÖ  Ú,r×  rŸ  r¡  r‚   )rþ   rí   r¤  rc   r¦   úOptional[BlockParameters])r’   rí   r¤  rc   r¦   ré  )r¦   z Optional[BlockDescriptorOptions])AÚprepare_indexingr�   r   r!  ÚoperatorÚ
attrgetterrŠ   r‹   rô   r   r¸   rð   r   rÄ   rÆ   rÇ   rÈ   rŸ  rû   rÊ   rË   rÌ   rÍ   rÎ   rÏ   rÐ   r   r�   r„  r"   rµ   r›  rü  rÒ   rÌ  Úvar_listrÓ   Úallow_block_ptrrÏ  r  Úis_indirect_indexingræ  r§  r“   rF  rÔ   r×   r  r<  re  rø   rY  rE  ÚcopyÚissubsetÚ
differencerW  r–  rX  r  rÖ   rî  r   r"  rÎ  )%r  rþ   rš  r›  rœ  rA  rB  Ú
index_varsr	  rÚ   rÛ   rÝ   Úprefix_matchesÚ
need_denseÚ
have_denseÚhave_loop_varsÚdense_mask_varsrá   rØ  r×  rü   rÿ   rù   rß  Ú
mask_shapeÚxyzrÚ	tmp_masksr  Úexpand_listr  rg  Úexpand_shape_strrW  rû   r¥  rÃ  r¼  s%   ```   `                          @@@@r†   rD  ÚTritonKernel.indexing£
  s‹  ÿ€ ð ×%Ñ% eÓ,ˆØ×'Ñ'ˆ
Øˆ
ä%/£\ˆ	Ü˜*¬(×*=Ò*=¸fÓ*EÕFˆCÜ˜c¤5§<¡<×0Ñ0Ð0Ð0Ø#÷ ¤~Ø”]×2Ñ2ó(ˆJö ÙÜ ¤T§X¡X×.Ñ.àŸ(™(×.Ñ.¨s¯x©xÑ8�Ø× Ñ  ×!2Ñ!2Ö3ÜØä×%Ñ%Ü—I‘IÜ×)Ñ)Ü—J‘JÜ—J‘JÜ×'Ñ'ð÷
ñ 
ò ô
 !.× 9Ò 9ó"â 9˜Ü% c×0ó %”J˜tÔ$Ù 9ð ð "ô
 �~Ó&¨!Ó+ØÜ˜>Ó*¨aÓ/ÐNÐ3CÀCÇHÁHÀ:Ð1NÓNÐ/Ø—‘ °Ñ!2Ð 3°4Ð8×9ñC GôH �M‰M×(Ñ(÷ +Ø÷+à�‰ dÐ*÷ð �q‰jð	 	ð ˆ
ØˆÜ+5«<ˆà×+Ñ+Ö-ˆDØ×&Ñ& t§}¡}×5Ñ5Ø!%‘à"�
Ø×Ñ 4§;¡; -¨tÐ 4Ö5ñ .ö ˜t×3×3¼¿¹×8S×8Sæ-Ø1×=Ñ=×?Ñ?ö "Ø—O—OÜ�I Ñ/Ó0°AÓ5Ø×-Ñ-¨e×4Ñ4Þà× Ñ  JÓ.ðØ!ðØ/Bðà*ôð,xØ!ðxØ/Bðxà*÷xðtØ ðà/ðð +÷ð ÷(_ô _ñD 'Ó(ˆGØÑ"Ø�Øˆ
Ø'+ˆØ×%Ñ% eÓ,ˆ	ö	Lô ! ×'Ò'ö œS ×!5Ñ!5Ó!7Ó8¸AÓ=Ù+:Ó+<Ñ(�
™Lä  ! ¤s¨4×+?Ñ+?Ó+AÓ'BÑ!BÓC�
Ü$ a S¬3¨t×/CÑ/CÓ/EÓ+FÑ%FÓG�à" : ,¨b°°¸;ÐGˆIØ× ×  D×$8×$8Ü&ô 'à $× 0Ò 0ó'ó ‘	ô '›L�	Ø��Ø—‘˜dŸo™oÔ.Ü"ØØØØØØ)ñð ÷ ŸjØ×$×$Ð$¨×)>×)>Ð)>ð. 'Ÿ^™^Ó-�
Ø—?—?Ø—N‘N 4§?¡?Ô3ä!Ò"HÓI�Ø$×-Ñ-¨d×3Ñ3Ø *× 5Ñ 5°dÓ ;�IØ#Ÿ-™-›/�CÜ% cÔ+<×=Ñ=Ð=Ð=Ø×&Ñ& sÔ+Ø×%Ñ% c§m¡mÔ4ð %×-Ñ-¨d×3Ó3ð  #˜e¤c¨$×*>Ñ*>Ó*@Ó&AÑA�Û&�DÜ% d¬C×0Ñ0Ð0Ð0Ø $× 7Ñ 7Ö 9˜ØŸ?™?¨4¯;©;×7Ó7Ø"&§/¡/˜CÜ#-¨c´3×#7Ñ#7Ð7Ð#7Ø/3×/CÑ/CÓ/EÀcÑ/J˜K¨Ó,ó	 !:ñ 'ð ! 3§8¡8¬C´°[Ó,AÓ#BÑBÀSÑH�
Ü$ [Ó1�Ø.¨y¨k¸¸J¸<ÀqÐI‘	á+:Ó+<Ñ(�
˜LØ.¨y¨k¸¸J¸<ÀqÐI�	Ø+‘	Þ¦JÙ-<Ó->Ñ*Ð˜lØ*¨9¨+°RÐ8HÐ7IÈÐKˆIØ'ˆIàÑÞžZÙ"1Ó"3‘�‘<à!�æÜ" M ?Ó3ˆIà�?�?Ø�M‰M˜$Ÿ/™/Ô*à×Ñ˜)Ô$äØØØØØØ%ñ
ð 	
ùòa"s   Æ ]1Æ]1c                óZ  • UR                  5       n[        U[        5      (       a  U(       a  U(       a
  US:X  d   eSnO'U(       d  SnOU(       a  US:X  d   eSU< S3nOSU< 3nU R                  (       a3  U R                  S   R
                  (       a  UR                  5       (       d  UR                  (       GaÝ  UR                  (       a!  X R                  ;   a  U R                  U   nXd4$ UR                  USS9nU R                  R                  U5      nU(       a  [        U5      U4$ [        U R                  5      n	[        U[        5      (       a  SU	 3nOS	U	 3nU R                  R!                  U["        R$                  / S
9n
U R                  R'                  Xz5        [)        X SU 35      nUR                  (       a*  U R*                  R-                  U5        X`R                  U'   OU R.                  R-                  U5        [        U[        5      (       au  XR0                  U'   [2        R4                   HS  nUR7                  U5      n[9        S U 5       5      (       a  M-  U R:                  U   nXn;  d   SU SU S35       eXÞU'   MU     Xd4$ UR                  U5      nXd4$ )a%  Generate a block pointer or tensor descriptor for Triton kernel operations.

This method creates either a block pointer (for regular Triton operations) or
a tensor descriptor (for TMA operations) based on the indexing type. It handles
caching and reuse of descriptors for performance optimization.

Args:
    name: The name of the buffer/tensor being accessed
    var: The variable name for the pointer
    indexing: Block pointer options or tensor descriptor options containing
             indexing information and boundary check settings
    other: Additional parameters string (e.g., padding options)

Returns:
    A tuple containing:
    - block_descriptor: The generated block pointer or tensor descriptor variable name
    - other: Modified additional parameters string with boundary check options
ú, other=0.0r°   ú, boundary_check=z, padding_option='zero'rK  F)rs  rA  Útma_descriptorr0  ú = c              3  óž   #   • U  HC  n[         R                  R                  R                  U[        R
                  " S 5      5      v •  ME     g7f©r   N)rL   rP  rH  rQ  r‹   r�   )rö   rº   s     r†   r  Ú1TritonKernel.codegen_block_ptr.<locals>.<genexpr>Þ  sC   é € ð ò +: ô ŸG™G×,Ñ,×DÑDØ &¬¯ª°aÓ(8÷ð ò +:ùs   ‚AAz#duplicate advancement for pointer 'z' at type 'r>  )r‰  rŠ   r»  rY  re  rƒ  r	  r7  r?  rÃ  rÆ   rf  rÔ   ÚnextrE  rÈ  Únamedvarrœ   Úuint64ri  rU   r@  rh  re  rG  r¸   rð   rÒ  rR  rH  )r  rÈ   rÚ   rD  rf  ÚcheckÚblock_descriptorÚblock_ptr_lineÚ	block_varÚblock_descriptor_idÚ	named_varÚ	line_bodyrÝ   Úadvance_offsetsÚadvancementss                  r†   Úcodegen_block_ptrÚTritonKernel.codegen_block_ptr‡  s®  € ð2 ×'Ñ'Ó)ˆÜ�hÔ 7×8Ñ8Þžð  Ó-Ð-Ð-Ø�øæà‘ÞØ Ó-Ð-Ð-Ø+¨E©9Ð4KÐL‘à+¨E©9Ð5�ð ×!×!Ø× Ñ  Ñ$×,×,Ø×#Ñ#×%Ñ%Ø××ÐØ× ×  S×,?Ñ,?Ó%?ð $(×#6Ñ#6°sÑ#;Ð ðh  Ð&Ð&ðe "*§¡°¸e Ð!D�Ø ŸH™H×,Ñ,¨^Ó<�	ö Ü˜y›>¨5Ð0Ð0ä&*¨4×+<Ñ+<Ó&=Ð#Ü˜h¬×8Ñ8Ø)2Ð3FÐ2GÐ'HÑ$à)7Ð8KÐ7LÐ'MÐ$Ø ŸH™H×-Ñ-Ø$¬E¯L©LÀð .ð �	ð —‘—‘˜^Ô7ä(¨Ð1CÀ3À~ÐFVÐ/WÓX�	Ø×$×$Ø—M‘M×+Ñ+¨IÔ6à/?×'Ñ'¨Ò,à—I‘I×'Ñ'¨	Ô2ä˜h¬×8Ñ8ð BF×,Ñ,Ð-=Ñ>ô !.× =Ô =˜Ø*2×*BÑ*BÀ4Ó*H˜ô ñ ñ +:ó	÷ ñ ñ %à'+×'@Ñ'@ÀÑ'F˜Ø/ÓCð ØAÐBRÐASÐS^Ð_cÐ^dÐdeÐfóÐCð :IÐ%5Ó6ñ! !>ð&  Ð&Ð&ð  (Ÿ™¨sÓ3ÐØÐ&Ð&rˆ   c                ó^  • SU SUR                    S3n[        [        UR                   UR                  5      5       HG  u  nu  px[        R
                  R                  R                  Xx5      (       d  M8  SUR                  U'   MI     UR                  UUR                   UR                  SSS9nU S[        [        R
                  R                  U5      5       S3n[        U[        5      (       a  SU SU U S3$ U S	[        R                  R!                  UR"                  5       SU S3$ )
NrŸ  r   r¡  FT©r¬  r­  r»  ú	tl.store(z.store()r3  r–  r¤  r0  rL   rP  rH  rQ  r2  r³  r;  r‡  Ú	get_dtyperŠ   rÈ  rÅ   r§  rB  )	r  rÈ   rD  rA  rª  rf  r…  rg  Úbroadcast_dims	            r†   Úcodegen_block_ptr_store_lineÚ)TritonKernel.codegen_block_ptr_store_lineï  s-  € ð # 5 '¨¨H×,@Ñ,@Ð+AÀÐCˆô *3Ü�×$Ñ$ h×&>Ñ&>Ó?ö*
Ñ%ˆCÑ%�#ô �w‰w×Ñ×7Ñ7¸×KÓKØ27�×*Ñ*¨3Ó/ñ	*
ð ×6Ñ6ØØ× Ñ Ø× Ñ Ø Øð 7ð 
ˆð �'˜Ô/´·±×0AÑ0AÀ$Ó0GÓHÐIÈÐKˆÜ�h¤×0Ñ0Ø˜y˜k¨¨E¨7°5°'¸Ð;Ð;Ø�˜G¤A§H¡H×$9Ñ$9¸(×:JÑ:JÓ$KÐ#LÈBÈuÈgÐUVÐWÐWrˆ   c                ó  • U(       d  U(       d  g [        U[        R                  5      (       d   eU R                  USS S9n[        U[        5      (       d   eUR
                  nUR                  5       (       a  UR                  OS nU(       a  [        U R                  U5      5      OS nU R                  Xc(       a  SOS X‡5      n	U R                  U5      n
U R                  R                  X©S[        R                  S9  g )NFr@  Ú0)Ú
assignmentrƒ  )rŠ   r‹   rŒ   rD  rø   rù   r  r#  ÚtexprÚrename_indexingÚindirect_assertÚget_load_bufferrÆ   r4  rœ   rŒ  )r  r’   rÞ  Úlowerrõ   rD  rù   r#  Úsize_strÚlineÚbuffers              r†   Úcheck_boundsÚTritonKernel.check_bounds  sÖ   € ö žØä˜$¤§
¡
×+Ñ+Ð+Ð+Ø—=‘= °ÐRV�=ÐWˆÜ˜(¤O×4Ñ4Ð4Ð4à×&Ñ&ˆ	Ø(0×(9Ñ(9×(;Ñ(;�8×$Ò$ÀˆÞ8=”5˜×-Ñ-¨dÓ3Ô4À4ˆð ×#Ñ#Ø�e‘s¨¨xó
ˆð ×%Ñ% hÓ/ˆØ�‰×Ñ˜&°5ÄÇÁÐÒLrˆ   c                ó&  • UR                  5       (       d  UR                  5       (       a  U R                  $ U R                  (       a?  U R                  S   R
                  (       a!  UR                  5       (       d  U R                  $ U R                  $ )NrK  )	r  r  r5  rY  re  rƒ  r	  re  Úloads)r  rD  s     r†   r"  ÚTritonKernel.get_load_buffer&  sk   € Ø× Ñ ×"Ñ" h×&:Ñ&:×&<Ñ&<à—<‘<Ðà×!×!Ø× Ñ  Ñ$×,×,Ø×'Ñ'×)Ñ)ð —9‘9Ðà—:‘:Ðrˆ   ztl.extra.cuda.gdc_wait()z%tl.extra.cuda.gdc_launch_dependents()c                 óÚ  • [         R                  R                  R                  R                  (       d  g[        [        R                  [         R                  R                  R                  5      (       a  g[         R                  R                  (       a  g[        R                  R                  5       R                  S:H  =(       a$    [         R                  R!                  5       S   S:¬  $ )NFrH  r   r  )rœ   r4  r"   rµ   Ú
enable_pdlrŠ   rL   rÅ   Úselect_algorithmÚTritonTemplateKernelrð  rñ  rP  r
  rA  rH  r  r‚   rˆ   r†   Ú_enable_pdl_codegenÚ TritonKernel._enable_pdl_codegen8  s‘   € ä�‰×%Ñ%×,Ñ,×7×7ØÜ”a—h‘h¤§¡× @Ñ @× UÑ U×VÑVØä�=‰=××Øä�G‰G×/Ñ/Ó1×6Ñ6¸&Ñ@÷ ;Ü—
‘
×0Ñ0Ó2°1Ñ5¸Ñ:ð	
rˆ   ©Úconsider_readsc               ó  ^^^^• U R                  5       (       d  g [        R                  R                  m[        R                  R
                  b$  [        R                  R
                  R                  OS mUUU4S jmU(       d   eTS L =(       d    [        U4S jU 5       5      nU(       d  g U R                  (       a*  U R                  S   R                  (       a  U R                  nUR                  U R                  5        g )Nc                óÌ   >^ • Tc   eTR                   R                  nT(       a+  [        R                  " UTR                   R                  5      n[        UU 4S jU 5       5      $ )Nc              3  óŠ   >#   • U  H8  nTTR                   R                  UR                  UR                  5      :H  v •  M:     g 7frš   )Úmutation_renamesrè  rÈ   )rö   ÚwÚcurrent_nodeÚdeps     €€r†   r  ÚOTritonKernel._handle_pdl_before_access.<locals>.matching_dep.<locals>.<genexpr>U  s8   øé € ð â"�Að �|×4Ñ4×8Ñ8¸¿¹ÀÇÁÓHÖHÚ"ùs   ƒA A)Úread_writesÚwritesr½  r¾  Úreadsrz   )r:  Ú	prev_depsr3  r9  Ú	prev_nodes   ` €€€r†   Úmatching_depÚ<TritonKernel._handle_pdl_before_access.<locals>.matching_depP  sY   ù€ ØÑ(Ð(Ð(Ø!×-Ñ-×4Ñ4ˆIÞÜ%ŸOšO¨I°y×7LÑ7L×7RÑ7RÓS�	Üõ á"óó ð rˆ   c              3  ó4   >#   • U  H  nT" U5      v •  M     g 7frš   r‚   )rö   r»  rA  s     €r†   r  Ú9TritonKernel._handle_pdl_before_access.<locals>.<genexpr>[  s   øé € Ð,SÂlÀ©\¸!¯_¨_Âlùrº  rK  )r0  rL   rÅ   r9  rP  Ú	schedulerÚprevious_noderz   rY  re  rƒ  re  rh  ÚGDC_WAIT)r  Úwait_bufferr3  ÚdependenciesÚ	need_waitr9  rA  r@  s     `  @@@r†   Ú_handle_pdl_before_accessÚ&TritonKernel._handle_pdl_before_accessF  s¶   û€ ð ×'Ñ'×)Ñ)ØÜ—x‘x×,Ñ,ˆä/0¯w©w×/@Ñ/@Ñ/LŒA�G‰G×Ñ×+Ò+ÐRVð 	÷	ö Ðˆ|Ø Ð%×S¬Ô,SÁlÓ,SÓ)Sˆ	ÞØà× ×  T×%5Ñ%5°bÑ%9×%A×%AØŸ)™)ˆKà×Ñ˜dŸm™mÕ,rˆ   c                ó4  • U R                  5       (       d  g UR                  S:”  a  g U R                  (       a*  U R                  S   R                  (       a  U R
                  nUR                  U R                  5        UR                  U R                  5        g )NrN   rK  )	r0  Ú	use_countrY  re  rƒ  rA  rh  rG  Ú
GDC_LAUNCH)r  Úlaunch_bufferÚ
result_vars      r†   Ú_handle_pdl_after_loadÚ#TritonKernel._handle_pdl_after_loadd  sr   € Ø×'Ñ'×)Ñ)ØØ×Ñ !Ó#Øà× ×  T×%5Ñ%5°bÑ%9×%A×%AØ ×2Ñ2ˆMð 	×Ñ §¡Ô.Ø×Ñ §¡Õ0rˆ   c                ó@  • / nSnS nUR                    H‚  n[        U5      [        L a  U R                  U;   a  U(       a  M.  Sn[        U5      [        L a/  U R                  U;   a  Ub  UR                  U5        [        U5      nUR                  U5        M„     X!l         g )NFT)Ú_linesrA  rÔ   rG  rO  rW  r�   rØ  )r  ÚcodeÚ	new_linesÚhas_waitÚprevious_launchÚls         r†   Ú_filter_pdlÚTritonKernel._filter_pdlt  s†   € Øˆ	ØˆØˆØ—”ˆAÜ�A‹wœ#Š~ $§-¡-°1Ó"4ÞÙà#�HÜ�A‹wœ#Š~ $§/¡/°QÓ"6Ø"Ñ.Ø—M‘M /Ô2Ü"% i£.�Ø×Ñ˜QÖñ ð  �rˆ   c                óN   • U R                   R                  [        XU5      5        g rš   )Úsaved_partial_accumulaterØ  rd   )r  rÈ   rƒ   rÃ  rq  s        r†   rr  ÚTritonKernel.partial_accumulate…  s"   € ð 	×%Ñ%×,Ñ,Ü˜d°CÓ8õ	
rˆ   c                óB  ^^^^^• U R                   R                  T5      nU R                  mTT==   S-  ss'   [        nU R	                  U5      mUn[
        R                  R                  T5      nU R                  USU R                  U USSS9S9n[        U[        5      (       a'  U R                  UR                  5      (       a  SU l        UR                  5       mUR!                  5       n[#        S U R%                  U5      R'                  5        5       5      n	U R)                  U5      (       a  Sn
OiU	(       d  Sn
O_U R*                  (       aL  U R,                  S   R.                  (       a.  UUUUU4S	 jnTT   mS
n
[0        R2                  " [4        SU5      nOSn
U(       d  T(       aB  UR7                  5       (       a-  U R8                  (       a  S[;        U R8                  5       3nOSnOSn Sn[<        R>                  R@                  (       a"  U RB                  RE                  5       nUT   S:„  n U R)                  U5      (       + =(       a(    U R*                  (       + =(       a    U(       + =(       a    U	nSnU(       a  SnSnSn[G        T5      (       al  UnU[H        RJ                  [H        RL                  4;   aB  [<        R>                  RN                  (       a  [H        RP                  nOUS[S        U5       S3-  nSnGOÝ[        U[T        [V        45      (       a›  U RY                  TX7U5      u  nn[        U[T        5      (       a  SU U U
 U S3nO/U S[
        RZ                  R]                  UR^                  5       S3nURa                  UURb                  URd                  SSS9nURd                  nO‹[g        U5      (       a  SU SU S3nURh                  nSnOcSU SURj                   SURl                   U
 U U S3
nURn                  (       a  URn                  nO[p        Rs                  UR                  5      nU[H        RJ                  [H        RL                  4;   a4  [<        R>                  RN                  (       a  US-  n[H        RP                  nU[H        Rt                  :X  a0  [H        Rv                  Rx                  c  US-  n[H        Rt                  nU R{                  U5      nU R}                  UT5        U R~                  R�                  UU" U5      UUS9nU Rƒ                  UU5        UR„                  S:”  a  TT==   S-  ss'   [        U[†        5      (       d   eURˆ                  UlD        U(       aÐ  SU SU S3nU R~                  R�                  UUXgRn                  S9nURˆ                  (       a‘  URŠ                  (       a  SnOU[H        Rt                  :X  a  S nOS!nU R8                  (       a  [;        U R8                  5      OUnS"URl                   SU SU S3nU R~                  R�                  UUUURŒ                  S9nU R*                  (       a  UR�                  5       (       d"  T(       d  U R�                  R“                  U5        U$ )#zS
Load from the memory location 'name', offset by some indexing expression 'index'.
rN   TF©r­  r  r@  c              3  ó*   #   • U  H	  oS :H  v •  M     g7fr›  r‚   )rö   rU  s     r†   r  Ú$TritonKernel.load.<locals>.<genexpr>´  s   é € ð 
ÚM�q�ŽFÒMùr\  z, eviction_policy='evict_last'rK  c                 ó6   >• TT   T :”  a  T(       d  T(       a  gg)NÚ
evict_lastÚevict_firstr‚   )Úexpected_countr	  Úindirect_indexingÚload_countsrÈ   s   €€€€€r†   Údecide_laterÚ'TritonKernel.load.<locals>.decide_later½  s   ø€ Ø˜tÑ$ ~Ó5ÞÖ"3à'Ø$rˆ   z, eviction_policy='<EP>'z<EP>r°   z, other=r   z, cache_modifier='.cg'Nr»  r¡  r‚   r_  z.load(r  r½  r@  r	  r²  rN  r0  rŸ  r   z0.0ÚTruer  r1  )JrÂ  rb  rK  r   rï  rL   rP  r  rD  Útma_compatibility_checker_clsrŠ   rø   rc  rþ   rU  r	  r  rz   Úget_strides_of_loadr¿  Úis_broadcastedrY  re  rƒ  r6  r7  r<   r  Ú_load_otherr`   r"   rµ   Úskip_l1_cacher[  Úbuffer_read_countsrj   rœ   r­  r®  r¬  r�   rF   rÈ  r»  r  rÅ   r§  rB  r³  r;  r3  r“   rü   rù   r#  rÿ   r¸   rã   r¥   rð  rñ  r"  rK  rÆ   r4  rR  rN  r–  rû   r‘  rÉ   r  rC  r„  )r  rÈ   rþ   rÚ   Ú	make_lineÚoriginal_indexrƒ  rD  r  Úis_coalescedÚeprj  rf  Úhas_read_depsrr  rq  ÚcachemodÚappend_broadcastrÉ   r%  r  Úload_bufferrQ  ÚzeroÚ	other_valrg  r	  rh  ri  s    `                       @@@@r†   rQ  ÚTritonKernel.loadŒ  s  ü€ ð �i‰i�o‰o˜dÓ#ˆØ×'Ñ'ˆØ�DÓ˜QÑÓÜCKˆ	Ø ×5Ñ5°eÓ<ÐØˆÜ—‘×!Ñ! $Ó'ˆØ—=‘=ØØØ&*×&HÑ&HØØØØð	 'Ið 'ð !ð 	
ˆô �h¤×0Ñ0°T×5NÑ5NØ�N‰N÷6
ñ 6
ð 26ˆDÔ.à×(Ñ(Ó*ˆ
Ø×*Ñ*Ó,ˆô ñ 
Ø ×4Ñ4°^ÓD×KÑKÔMó
ó 
ˆð ×Ñ˜~×.Ñ.Ø1‰BÞØ1‰BØ×"×" t×'7Ñ'7¸Ñ';×'C×'C÷%ñ %ð )¨Ñ.ˆNØ+ˆBÜ!×)Ò)Ô*:¸FÀLÓQ‰IàˆBæž:¨8×+<Ñ+<×+>Ñ+>Ø××Ø"¤=°×1AÑ1AÓ#BÐ"CÐD‘à%‘àˆEð	ð ˆÜ�=‰=×&×&Ø!%§¡×!AÑ!AÓ!CÐØ.¨tÑ4°qÑ8ˆMð	ð ×#Ñ# NÓ3Ô3÷ Ø×)Ñ)Ô)÷à!Ô!÷ð ð	 	ð ˆÞØ/ˆHàÐØ $ˆä# D×)Ñ)ØˆDð œŸ™¬¯©Ð7Ó7Ü—=‘=×7×7Ü!ŸM™M‘Eà˜d¤;¨uÓ#5Ð"6°aÐ8Ñ8�DØŠEô ˜(¤_Ô6MÐ$N×OÑOØ*.×*@Ñ*@Ø˜#¨ó+Ñ'Ð  %ô ˜h¬×8Ñ8Ø%Ð&6Ð%7¸°w¸r¸dÀ8À*ÈAÐN‘Dà.Ð/¨v´a·h±h×6KÑ6KÈH×L\ÑL\Ó6]Ð5^Ð^_Ð`�DØ×=Ñ=ØØ×(Ñ(Ø×(Ñ(Ø#'Ø#ð >ð �ð !×,Ñ,‘Ü& ~×6Ñ6Ø! #  d¨>Ð*:¸"Ð=�Ø#+×#6Ñ#6Ð Ø‘à! #  d¨8×+=Ñ+=Ð*>¸cÀ(×BSÑBSÐATÐUWÐTXÐY^ÐX_Ð`hÐ_iÐijÐk�ð ×(×(Ø$×1Ñ1‘Eä)×9Ñ9¸(¿.¹.ÓI�Eð œ%Ÿ-™-¬¯©Ð8Ó8Ü—M‘M×8×8àÐ)Ñ)�ÜŸ™�ØœŸ
™
Ó"¤u§}¡}×'8Ñ'8Ñ'@ð ˜Ñ&�ÜŸ
™
�à×*Ñ*¨8Ó4ˆØ×&Ñ& {°DÔ9Ø—X‘X×&Ñ&Ø™ 4›°¸Uð 'ð 
ˆ
ð 	×#Ñ# K°Ô<Ø×Ñ !Ó#Ø˜Ó Ñ"ÓÜ˜*Ô&7×8Ñ8Ð8Ð8Ø'×1Ñ1ˆ
ÔæØ% j \°Ð4DÐ3EÀQÐGˆDØŸ™×*Ñ*Ø˜T¨×6KÑ6Kð +ð ˆJð ×!×!Ø×*×*Ø ‘DØœeŸj™jÓ(Ø!‘Dà�Dà7;×7G×7G”M $×"2Ñ"2Ô3ÈTð ð # 8×#4Ñ#4Ð"5°R¸
°|À2ÀiÀ[ÐPQÐR�Ø!ŸX™X×.Ñ.Ø ¨U¸*×:JÑ:Jð /ð �
ð ×$×$¨X×-?Ñ-?×-AÑ-AÎ*Ø×"Ñ"×&Ñ& zÔ2àÐrˆ   c           	     óf  • U R                   R                  U5      nUn[        R                  R	                  U5      nSnUb  US:X  a  US:H  n	U R                  U USU	S9nU R                  USUSL US9n
[        U
[        5      (       a;  U R                  U
R                  5      (       a  U R                  R                  U5        XR                   R                  ;   nU R                  U5      nU(       a,  U(       a%  U R                  R!                  [#        US5      5        [        U
[$        [&        45      (       a)  U R)                  XU
5      u  pÞU R+                  XXÓU5      nGOsUc¿  U
R,                  n[/        U5      (       a‰  U R0                  (       a  USU R3                  5        S3-  nO`UR4                  bS  [7        S	 UR4                   5       5      (       d2  S
R9                  [;        [<        UR4                  5      5      nUSU S3-  nSU SU SU S
U
R>                   S3	nO±US:X  a�  SU l         U
R,                  n[/        U5      (       a`  UR4                  bS  [7        S UR4                   5       5      (       d2  S
R9                  [;        [<        UR4                  5      5      nUSU S3-  nSU SU SU S
U
R>                   S3	nO[C        SU 35      e[D        RF                  " 5       nU RH                  (       d;  U RJ                  (       a*  URM                  U RO                  XR                  5      5        U RQ                  U R                  USS9  U R                  R!                  [#        X5      5        U RH                  (       d  U RR                  RU                  U5        URW                  5         g)z^
store the 'value' to the memory location 'name', offset by some indexing expression 'index'.
NÚtmaTra  )r›  rA  rB  ztl.debug_barrier()ú.broadcast_to(r¡  c              3  ó>   #   • U  H  n[        U5      S :H  v •  M     g7f©rÃ   NrÜ  r:  s     r†   r  Ú%TritonKernel.store.<locals>.<genexpr>„  s   é € ð 9Ú+6 a”C˜“F˜c–Mª;ùó   ‚r   r  r½  r	  Ú
atomic_addc              3  ó>   #   • U  H  n[        U5      S :H  v •  M     g7fr‚  rÜ  r:  s     r†   r  rƒ  �  ó   é € Ð?²;¨aœC ›F cžM²;ùr„  ztl.atomic_add(z, sem='relaxed')zstore mode=r2  ),rÂ  rO  rL   rP  r  rm  rD  rŠ   rø   rc  rþ   rV  rØ  Úinplace_buffersro  Ústoresrh  rU   rÈ  r»  r  r  rù   r“   r<  rÞ  rÉ   rR  r   r"  rÔ   r#  Úatomic_add_foundr<  Ú
contextlibÚ	ExitStackrY  rQ  Úenter_contextÚguard_cooperative_storerK  rC  r„  Úclose)r  rÈ   rþ   rª  ÚmoderÚ   rt  rƒ  rB  r  rD  Ú
is_inplacero  r  rf  r%  Úindexing_strÚvalue_shapeÚ
exit_stacks                      r†   ÚstoreÚTritonKernel.storeF  sW  € ð �i‰i×Ñ˜tÓ$ˆØˆÜ—‘×!Ñ! $Ó'ˆà$(Ð!Ø‰<˜4 5›=Ø˜E‘MˆEØ(,×(JÑ(JØØØØð	 )Kð )Ð%ð —=‘=ØØØ˜d�lØ&?ð	 !ð 
ˆô �h¤×0Ñ0°T×5NÑ5NØ�N‰N÷6
ñ 6
ð ×,Ñ,×3Ñ3°DÔ9ð ŸY™Y×6Ñ6Ñ6ˆ
Ø×,Ñ,¨^Ó<ˆÞž.Ø�K‰K×!Ñ!¤,¨tÐ5IÓ"JÔKä�h¤Ô2IÐ J×KÑKØ&*×&<Ñ&<¸TÈÓ&QÑ#Ðà×4Ñ4ØÐ 0¸óŠDð ‰\ð $×-Ñ-ˆLÜ$ U×+Ñ+Ø×'×'à  n°T×5HÑ5HÓ5JÐ4KÈ1Ð$MÑM‘LØ—[‘[Ñ,´Sñ 9Ø+0¯;ª;ó9÷ 6ñ 6ð #'§)¡)¬C´°U·[±[Ó,AÓ"B�KØ  n°[°MÀÐ$CÑC�LØ˜s˜e 4¨ ~°S¸¸¸rÀ(×BSÑBSÐATÐTUÐV‰DØ�\Ó!Ø$(ˆDÔ!Ø#×-Ñ-ˆLä% e×,Ñ,Ø—K‘KÑ+ÜÑ?°5·;²;Ó?×?Ñ?à"Ÿi™i¬¬C°·±Ó(=Ó>�Ø .°°¸QÐ ?Ñ?�Ø# C 5¨¨\¨N¸#¸e¸WÀBÀx×GXÑGXÐFYÐYiÐj‰Dä%¨°D°6Ð&:Ó;Ð;ä×)Ò)Ó+ˆ
Ø×$×$¨×)C×)CØ×$Ñ$ T×%AÑ%AÀ$ÏÉÓ%TÔUà×&Ñ& t§{¡{°DÈÐ&ÑNØ�‰×Ñœl¨4Ó6Ô7à×$×$Ø×"Ñ"×&Ñ& uÔ-à×ÑÕrˆ   c                óZ   • U R                   R                  SU S[        U5       S35        g )Nztl.device_assert(r   r¡  )r5  rh  Úrepr)r  ÚcondÚmsgs      r†   Údevice_assert_asyncÚ TritonKernel.device_assert_async¤  s(   € Ø�‰×ÑÐ!2°4°&¸¼4À»9¸+ÀQÐGÕHrˆ   c                ó”   • U R                   R                  5       nUR                  [        USU S35      5        UR	                  5       $ )z“
For cooperative reductions only one thread block should write out the result.
We rotate which thread block does each write for better parallelism
zif rsplit_id == (z % RSPLIT):)ry  rî  rh  rU   Úindent)r  rÈ   r&  r…  s       r†   rŽ  Ú$TritonKernel.guard_cooperative_store§  sC   € ð
 ×8Ñ8×NÑNÓPˆØ×Ñœ dÐ.?À¸uÀKÐ,PÓQÔRØ�}‰}‹Ðrˆ   c                óŠ   • S nU H:  nUc  M  [        US5      (       d  M  Uc  UR                  nM,  X#R                  -  nM<     U$ )Nrû   )r³   rû   )r  Ú	variablesÚmasksÚelems       r†   Ú_combine_masksÚTritonKernel._combine_masks°  sG   € ØˆÛˆDØ‰|ÙÜ�t˜[×)Ó)Ø‘=Ø ŸN™N’Eà!§N¡NÑ2’Eñ ð ˆrˆ   c                óp  • U R                   R                  [        R                  5        U R                  R                  US   5      nU R                  US   5      n	U R                  US   5      n
U R                  US   5      nU(       a  U R                  R                  US   5      OSnU(       a  U R                  US   5      OSnU[        R                  :X  a  SnO"U[        R                  :X  a  SnO[        S5      eU R                  " U R                  US   /U(       a  US   /O/ Q76   U R                  R                  U R                  S	U S
U S
U	 S
U
 S
U S
U S
U S
U S
U S
U S
U S3UUR                  S9nU R!                  U R                  U5        U R#                  XU5      nUUl        U$ )z#
See [Note: Inductor bucketize op]
r   rN   r!   r   r§   rŸ  ztl.int64z5Bucketize only supports indexing with int32 and int64z'triton_helpers.bucketize_binary_search(r   z, )r0  )rN  r„  r/   ÚONE_ELEMENT_PER_THREADrÂ  rb  r§  rœ   rŒ  rF  r<  rK  r5  rÆ   r4  rÉ   rR  r¤  rû   )r  r¿  Ú
boundariesÚboundary_indicesÚindexing_dtyper  ÚsorterÚsorter_indicesÚboundaries_ptrÚboundary_sizeÚboundaries_underlying_numelÚboundary_strideÚ
sorter_ptrÚsorter_strideÚtriton_dtyperk  r¢  s                    r†   Ú	bucketizeÚTritonKernel.bucketize¼  sÂ  € ð$ 	×Ñ×Ñ¤× CÑ CÔDàŸ™Ÿ™¨°A©Ó7ˆØ×)Ñ)¨*°Q©-Ó8ˆØ&*×&7Ñ&7¸
À1¹Ó&FÐ#Ø×+Ñ+¨J°q©MÓ:ˆÞ39�T—Y‘Y—_‘_ V¨A¡YÔ/¸vˆ
Þ8>˜×)Ñ)¨&°©)Ô4ÀFˆàœUŸ[™[Ó(Ø%‰LØœuŸ{™{Ó*Ø%‰Lä%ØGóð ð 	×&Ò&Ø�L‰L˜* Q™-ð	
Þ:@¨F°1©I©;Àbó	
ð —‘×"Ñ"Ø�L‰LØ5°f°X¸RØÐ˜b  ¨rÐ2MÐ1NÈbÐQ`ÐPaÐacØÐ  Øˆn˜BØˆg�RØˆl˜"˜]˜O¨2ØÐð ðð !Ø—,‘,ð #ð 
ˆð 	×#Ñ# D§L¡L°&Ô9à×#Ñ# F¸nÓMˆØ ˆÔàˆrˆ   c                óž   • U R                  5       nUS:X  a  SU S3$ U R                  nS/X#-
  -  S/U-  -   nU SSR                  U5       S3$ ©	NrN   z!triton_helpers.promote_to_tensor(r¡  rÕ  r§   rÖ  r   r×  ©rÑ   rX  r   )r  rª  ÚndimsÚnreduceÚsizess        r†   Úreduction_resizeÚTritonKernel.reduction_resize÷  sh   € Ø×'Ñ'Ó)ˆØ�A‹:Ø6°u°g¸QÐ?Ð?à×)Ñ)ˆØ�˜™Ñ)¨V¨H°wÑ,>Ñ>ˆØ�˜˜$Ÿ)™) EÓ*Ð+¨1Ð-Ð-rˆ   c                óÐ   • U R                  5       nUS:X  a  SU S3U4$ U R                  nS/X4-
  -  S/U-  -   nUb  / US X4-
   QS/U-  Q7OS nU SSR                  U5       S3U4$ r·  r¸  )r  rª  rÉ   r¹  rº  r»  rÙ  s          r†   Úreduction_resize_and_shapeÚ'TritonKernel.reduction_resize_and_shape   sŸ   € Ø×'Ñ'Ó)ˆØ�A‹:Ø6°u°g¸QÐ?ÀÐFÐFà×)Ñ)ˆØ�˜™Ñ)¨V¨H°wÑ,>Ñ>ˆà=BÑ=NÐ9ˆeÐ'�u‘Ð(Ð9¨A¨3°©=Ò9ÐTXð 	ð �˜˜$Ÿ)™) EÓ*Ð+¨1Ð-¨yÐ8Ð8rˆ   c                ó  • U R                   S:X  a  U$ U R                  5       U R                   -
  nU R                  5       nUSU S/-   nU R                  R	                  U[        [        U5      XV5      U[        U5      S9$ )z3
Reshape to RBLOCK, collapsing all reduction dims.
rN   Nr?  r0  )rX  rÑ   rF  rÆ   r4  r¥  rÔ   r×   )r  r&  rª  rƒ  Útarget_ndimr«  Útarget_shapes          r†   Úreduction_collapse_dimsÚ$TritonKernel.reduction_collapse_dims  s‰   € ð ×"Ñ" aÓ'ØˆLà×-Ñ-Ó/°$×2IÑ2IÑIˆØ×,Ñ,Ó.ˆØ$ \ kÐ2°h°ZÑ?ˆØ�x‰x× Ñ ØÜœ3˜u›: }ÓCØÜ˜Ó%ð	 !ð 
ð 	
rˆ   c                óJ  ^ ^^^^:^;^<^=^>^?^@^A• SES jn[         R                  " U5       Vs/ s H  ofR                  PM     nn[         R                  " XT5      n[	        S U 5       5      (       aJ  [
        R                  " T[
        R                  5      m[
        R                  " T[
        R                  5      mT R                  (       d   e[        S T R                   5       5      nT R                  U5        [        U5      nT R                  (       a  UR                  T R                  5        T R                  S   R                  S   n	T R                   (       ab  T R#                  5       n
[%        U
5      S:¼  d   eU
 Vs/ s H  nSU;   d  SU;   d  M  UPM     nnS	S
R'                  U5       S3m;[)        U5      m@O)T R+                  5       m;[)        T R#                  5       5      m@T R-                  U;U U@4S jU5      nSnTS;   a  [/        U[(        5      (       a  Uu  pMT R1                  5       T R2                  -
  m<      SFU<UUU 4S jjm=        SGU=4S jjnU<UU>U 4S jnTTU4nUT R4                  R6                  ;   a  T R4                  R6                  U   $ [9        T5      n[;        T5      n[=        T R#                  5       5      nSUT<'   T R4                  R?                  U[)        U5      S9n[        S U 5       5      Ul         SR'                  U5      m:U:4S jmAT RB                  (       Ga”  [D        RF                  RI                  TT5      nS m?SHU UU?UA4S jjnTS:X  a  SnO|[/        U[(        5      (       a)  [K        UU5       VVs/ s H  u  nnU" UU5      PM     nnnO>TS:X  a/  T R4                  RM                  T RN                  XDR                  S9nO	U" UU5      nTS;   aÄ  [/        U[P        5      (       d   e[R        RT                  RW                  5       nU(       a"  S[Y        U5       ST R[                  U5       S3nO@[Y        T R4                  RM                  T RN                  S U	 S!U S"3UUR\                  S95      nS#S$S.T   m>U" T RN                  UUU5        UUl        GO¹TS%:X  a=  T R^                  (       a  T Ra                  UTUTAUT5      nGOŠT Rc                  TU5      nGOvTS&:X  aS  [/        U[d        5      (       d   eUu  nnn[)        UU 4S' jT Rg                  T RN                  UUUT<T5       5       5      nGOTS:X  a  T Ri                  TU5      nGO[/        U[P        5      (       d   eT=" T RN                  UUR                  5      u  nn n!T R4                  RM                  T RN                  UU U!S9nGO¢T R4                  Rk                  S(U 3U[)        T R#                  5       5      S9n"[D        RF                  Rm                  TT5      nT R-                  [n        U5      n[/        U[(        5      (       dÉ  TS:X  aŽ  T R#                  5       n
[%        U
5      S:¼  d   eU
 Vs/ s H  nSU;   d  SU;   d  M  UPM     nn[)        U5      U"l.        S	S
R'                  U5       S3m;T Rp                  Rs                  U" S)T; S
U S
U S35        O5T Rp                  Rs                  U" S)T R+                  5        S
U S
U S35        TS;   Ga  S(U S*3nT Rt                  Rw                  5       n#T Rp                  Rs                  U S)T R+                  5        S
[
        Rx                  " U#5      Rz                   S
T R[                  U#5       S35        S#S$S.T   m>Ub!  S[Y        U5       ST R[                  U#5       S3OU	 S+3n$T RN                  R}                  S,U" S-U S.T> S/U" S
U S
U S
U$ S0U" S1TA" U" S23U"5       S3U S1TA" U S23U5       S335        U" T R~                  UU"U5        GO7[�        T5      (       a  T Ra                  UTUTAUT5      nGOTS:X  Ga0  S(U S43n%S(U S53n&T Rp                  Rs                  U% S)T R+                  5        S6U S35        T Rp                  Rs                  U& S7T R+                  5        S
U S35        T RN                  R}                  S8U% S-U& S9U% S
U& S
U S
[‚        R„                   S:35        T RN                  R}                  S8U% S1TA" U% S23U%5       S8U& S1TA" U& S23U&5       S83	5        Un'T R4                  R?                  TU'R\                  S9n(T R‡                  T R~                  U'U(U%U&T<T5      nOØ[D        Rˆ                  " TT5      n)U)" U"U5      n*TS:X  a!  T RN                  Rs                  U" S1U* 35        O'T RN                  Rs                  U" S1TA" U*U"5       35        T[
        RŠ                  :X  aA  T R4                  RM                  T R~                  U" S;3[
        RŒ                  U"R\                  S9n"U" T R~                  UU"S5        T R^                  (       Ga–  [D        RF                  Rm                  TT5      n[Ž        R�                  " 5       n+T R~                  T R’                  4 H3  n,U,Rs                  S<5        U+R•                  U,R—                  5       5        M5     TS;   aª  T R~                  Rs                  U S=T R™                  U S>35       35        T R›                  U S?3TU5      n-T Rt                  Rw                  5       n#T R›                  UU#[
        Rx                  " U#5      Rz                  5      n.U" T R’                  UU-U.5        GOR[�        T5      (       a�  TS%:X  d   eUu  n/n0n1T R›                  U/[;        T5      US   5      n2T R›                  U0[;        T5      US@   5      n3T R›                  U1[;        T5      USA   5      n4T R�                  T R’                  U/U0U1U2U3U4T<T5	        O³TS:X  a|  Uu  n'n([/        U[d        5      (       d   eT R›                  U'[;        T5      US   5      n5T R›                  U([;        T5      US@   5      n6T R‡                  T R’                  U'U(U5U6T<T5        O1T R›                  U[;        T5      U5      n7U" T R’                  UU7S5        U+RŸ                  5         UT R4                  R6                  U'   [/        U[(        5      (       aÏ  [¡        SB U 5       5      (       d   eT R¢                  R¥                  U5        TSC;   a  [%        U5      S@:X  d   e[%        U5      U-  n[%        U5      [%        U5      :X  d   e[K        UU5       HJ  u  n8n9U9c   eU8R                  U9:w  d  M  T R~                  Rs                  U8 S1U8 SD[§        U95       S35        ML     U$ [/        U[¨        5      (       d   eT R¢                  R«                  U5        UR                  US   :w  a8  US   c   eT R~                  Rs                  U S1U SD[§        US   5       S35        U$ s  snf s  snf s  snnf s  snf )IzC
codegen reduction of value to Triton according the reduction_type
c                ó¬   • U R                   [        R                  [        R                  4;   a%  [        R
                  " U [        R                  5      $ U $ rš   )rƒ  rœ   r­  r®  r  rÞ  r�   rR  s    r†   Úmaybe_upcastÚ,TritonKernel.reduction.<locals>.maybe_upcast+  sF   € ð —;‘;ä—M‘MÜ—N‘Nðóô —’˜U¤E§M¡MÓ2ðð ðrˆ   c              3  óf   #   • U  H'  o[         R                  [         R                  4;   v •  M)     g 7frš   )rœ   r­  r®  r:  s     r†   r  Ú)TritonKernel.reduction.<locals>.<genexpr>;  s    é € ÐMº_¸”U—]‘]¤E§N¡NÐ3Ö3º_ùs   ‚/1c              3  ó>   #   • U  H  oR                    S 3v •  M     g7frâ  ©rÓ   r„  s     r†   r  rË  A  ó   é € ÐMÒ<L°DŸk™k˜]¨$Õ/Ò<Lùr„  rK  r   r   ÚXÚYrÖ  r   r×  c                ór   >• TR                   R                  TR                  SU  ST S3U R                  TS9$ )NrŸ  r   r¡  r0  )rÆ   r4  r5  rƒ  )r+  rÞ  r  r“  s    €€€r†   Ú<lambda>Ú(TritonKernel.reduction.<locals>.<lambda>a  s=   ø€ �d—h‘h×'Ñ'Ø—‘Ø" 1 # R¨Ð'7°qÐ9Ø—g‘gØ!ð	 (ñ rˆ   N)ÚargminÚargmaxc                ó¬  >• [        T	5      nT
R                  XT5      nT	S:X  a_  [        T
R                  5       5      S:H  nUR                  c   eU(       a  U S3nS/UR                  QSPnO=U S3n/ UR                  QSPnO'T
R                  U SU ST S3UR                  5      u  pVUb  U S	T
R                  U5       S3nOUR                  nXRU4$ )
z3
Helper to generate a reduction call, e.g. tl.sum.
r/  r‰  z[None,:,:,None]rN   z
[:,:,None]r)  r   r¡  r»  )r‡   rÄ  r�   rF  rÉ   r¿  rl  rƒ  )r&  rª  Úresult_typeÚtriton_reduction_fnÚis_bmmrk  rÉ   rg  rƒ  rƒ   r  s          €€€€r†   Úfinal_reductionÚ/TritonKernel.reduction.<locals>.final_reductionr  s   ø€ ô #@ÀÓ"OÐà×0Ñ0°ÀÓFˆEØ Ó&ä˜T×1Ñ1Ó3Ó4¸Ñ9�Ø—{‘{Ñ.Ð.Ð.ÞØ %˜w oÐ6�FØÐ0 §¡Ð0¨aÐ0‘Eà %˜w jÐ1�FØ-˜eŸk™kÐ-¨1Ð-‘Eà $× ?Ñ ?Ø*Ð+¨1¨U¨G°2°c°U¸!Ð<¸e¿k¹kó!‘�ð Ñ&Ø"˜8 4¨×(9Ñ(9¸+Ó(FÐ'GÀqÐI‘à#Ÿk™k�à¨Ð-Ð-rˆ   c                óL   >• T" XU5      u  n  nU R                  U SU 35        g)z=
Generate a reduction and assign it to an existing variable.
r  N)rz  )r&  rQ  rª  r×  rW  rÚ  s        €r†   Úfinal_reduction_defineÚ6TritonKernel.reduction.<locals>.final_reduction_define“  s.   ø€ ñ *¨&¸ÓE‰KˆE�1�aØ�M‰M˜Z˜L¨¨E¨7Ð3Õ4rˆ   c                óÆ   >• TR                  XT5      nTR                  XT5      nU R                  SU SU ST SU SU ST SU STR                  U S35       S	35        g )
Nú                z_val, z_idx = triton_helpers.z_with_index(r   ú)
                r  Ú_idxú
                )rÄ  rz  r¼  )r&  rQ  rª  rþ   rg  rƒ  Úroot_opr  s       €€€€r†   Úfinal_argreduceÚ/TritonKernel.reduction.<locals>.final_argreduce   s“   ø€ Ø×0Ñ0°ÀÓFˆEØ×0Ñ0°ÀÓFˆEØ�M‰MðØ�˜F : ,Ð.DÀWÀIÈ\ÐZ_ÐY`Ð`bÐchÐbiÐikÐloÐkpð qØ�˜C × 5Ñ 5¸¸ÀDÐ6IÓ JÐKð Lðõrˆ   rÃ   r0  c              3  óP   #   • U  H  n[        US    5      (       a  M  Uv •  M     g7fr  )rB   )rö   rÚ   s     r†   r  rË  µ  s!   é € ð *
Ú �CÔ(;¸CÀ¹F×(C�C‰Cš5ùs   ‚&�	&r  c                óB   >• T(       d  U $ [         R                  TX5      $ rš   )r3  r#  )ÚtvalÚfvalr™  s     €r†   r9  Ú*TritonKernel.reduction.<locals>.where_condº  s   ø€ ÞØ�Ü(×.Ñ.¨t°TÓ@Ð@rˆ   c                ó~  • UR                   (       d  U[        R                  :X  a  U $ X:X  d  U S:X  a  U $ U [        R                  " U5      R                  :X  a   [        R                  " U5      R                  $ U [        R                  " U5      R
                  :X  a   [        R                  " U5      R
                  $ U $ )z7update reduction constant mask value to match dst_dtyper   )r‘  rœ   r¥   Úiinfor{   r|   )rî  rÕ  rÖ  s      r†   Úupdate_constant_dtypeÚ5TritonKernel.reduction.<locals>.update_constant_dtypeÂ  s‹   € ð
 ×.×.°)¼u¿z¹zÓ2IØ#�OàÓ)¨X¸«]Ø#�OàœuŸ{š{¨9Ó5×9Ñ9Ó9Ü Ÿ;š; yÓ1×5Ñ5Ð5Ø¤§¢¨YÓ!7×!;Ñ!;Ó;Ü Ÿ;š; yÓ1×5Ñ5Ð5à#�Orˆ   c                óØ   >• T" UTU R                   5      nTR                  [        U5      nTR                  R	                  TR
                  T" X5      U R                   U R                  S9$ )Nr0  )rƒ  Ú_map_tuple_or_scalarr`   rÆ   r4  r5  rÉ   )rª  r6  Údefault_strr  rÕ  rî  r9  s      €€€€r†   Ú_mask_valueÚ+TritonKernel.reduction.<locals>._mask_valueÔ  s`   ø€ Ù/°¸ÀEÇKÁKÓP�Ø"×7Ñ7¼ÀwÓO�à—x‘x×(Ñ(Ø—L‘LÙ˜uÓ2ØŸ+™+ØŸ+™+ð	 )ð ð rˆ   Úonline_softmax_reducer/  r‚  )rÕ  rÔ  r)  rã  r¡  rŸ  zindex, z.shape)r{   r|   Úwelford_reduceÚwelford_combinec              3  ót   >#   • U  H-  u  pTR                   R                  TR                  UTUS 9v •  M/     g7f)r0  N)rÆ   r4  r5  )rö   rª  rÉ   rƒ  r  s      €€r†   r  rË    s:   øé € ð #ò)™˜ð —H‘H×%Ñ% d§l¡l°EÀÈeÐ%ÕTò)ùs   ƒ58rW  ú = tl.full(Ú_indexrþ   rà  ú_next, z_next = triton_helpers.z%imum_with_index(
                    ú#
                )
                r  Ú_nextrã  Ú_maxÚ_sumz, float('-inf'), ú = tl.zeros(z
                    zG_next = triton_helpers.online_softmax_combine(
                        z+
                    )
                    z.to(tl.int8)zif HAS_RSPLIT:z_bval = Ú_valÚ_bvalrN   r!   c              3  óB   #   • U  H  n[        U[        5      v •  M     g 7frš   )rŠ   r–  r:  s     r†   r  rË  ï  s   é € ÐLÂ¸A”z !Ô%6×7Ð7Âùs   ‚)rö  rõ  r»  )rª  rT   r¦   rT   )rª  rT   r×  r   r¦   z1tuple[str, Optional[torch.dtype], BlockShapeType])rQ  rT   rª  rT   r×  r   r¦   r§   ©r¦   rT   )VÚpytreeÚtree_leavesrƒ  Útree_maprz   rœ   rI  r�   rY  r   re  rÎ  r!  rü  rØ  rÓ   rE  rF  r�   r   r×   rÞ  rñ  rŠ   rÑ   rX  rÆ   Úreduction_cacher�  r�  ra  rg  rû   r  r#   r1  r2  r¤  r4  r5  rT   rL   rÅ   rE  rÔ   rl  rÉ   rQ  rö  Úwelford_reduce_fallbackr   Ú_welfordÚ prepare_softmax_twopass_fallbackr  Údefault_accumulatorr`   re  rh  r[  Úselect_index_dtyperí  r{   rz  rA  r@   r"   r  Ú%online_softmax_reduce_final_reductionÚget_reduction_combine_fnr¥   r†  r‹  rŒ  rB  r�  rž  r¼  Ú*codegen_cooperative_reduction_peer_combineÚwelford_reduce_final_reductionr�  rR  rC  rŸ  r„  r–  r„  )Br  rƒ  rÕ  rƒ   rª  rÈ  rÃ  Úoriginal_dtypesr¢  Úreduction_range_prefixÚdense_sizesrÞ  Úxy_sizes_onlyÚlogical_indexrÝ  rå  rj  Úacc_typeÚtorch_acc_typeÚresult_shaperQ  r6  ró  Úmasked_valuer+  r»  Úaccumulator_dtypeÚaccumulator_indexÚmeanÚm2ÚweightÚ_resultÚ_dtypeÚ_shapeÚaccumulatorræ  rK  Úaccumulator_maxÚaccumulator_sumÚ
result_maxÚ
result_sumÚ
combine_fnÚupdatedr”  ÚbufÚpeer_valÚpeer_idxÚresult_meanÚ	result_m2Úresult_weightÚ	peer_meanÚpeer_m2Úpeer_weightÚpeer_maxÚpeer_sumÚpeersrÚ   Ú
orig_dtyper™  rÞ  rg  rÚ  rä  rî  r“  r9  sB   ````                                                      @@@@@@@@r†   Ú	reductionÚTritonKernel.reduction   s\  ÿû€ ô	ô 17×0BÒ0BÀ5Ô0IÓJÒ0I¨Ÿ9œ9Ñ0IˆÐJÜ—’ Ó4ˆÜÑM¹_ÓM×MÑMä×+Ò+¨I´u·}±}ÓEˆIÜ×'Ò'¨¬u¯}©}Ó=ˆEà×$×$Ð$Ð$ÜÑM¸D×<LÒ<LÓMÓMˆØ×Ñ˜%Ô Ü�u“ˆØ�?�?Ø�L‰L˜Ÿ™Ô)Ø!%×!1Ñ!1°"Ñ!5×!<Ñ!<¸QÑ!?Ðð × × Ø×.Ñ.Ó0ˆKÜ�{Ó# qÓ(Ð(Ð(Ù.9ÓXªk d¸SÀD»[ÈCÐSWÉKŸT©kˆMÐXØ  §¡¨=Ó!9Ð :¸!Ð<ˆNÜ Ó.‰Kà!×0Ñ0Ó2ˆNÜ × 4Ñ 4Ó 6Ó7ˆKð ×)Ñ)öð ó
ˆð ˆØÐ1Ó1Ü˜%¤×'Ñ'Ø',Ñ$�à×%Ñ%Ó'¨$×*AÑ*AÑAˆð	.àð	.ð /ð	.ð ?÷		.ò 	.ðB	5à#ð	5ð ð	5ð /ð		5ð
 ÷	5÷	ð 	ð  °Ð6ˆ	Ø˜Ÿ™×0Ñ0Ó0Ø—8‘8×+Ñ+¨IÑ6Ð6ä" 9Ó-ˆÜ)¨)Ó4ˆÜ˜D×0Ñ0Ó2Ó3ˆØˆ�SÑØŸ(™(Ÿ/™/Ø ¬¨lÓ(;ð *ð 
ˆ
ô  *ñ *
Ù ó*
ó  
ˆ
Ôð �z‰z˜%Ó ˆõ	Að
 ×$×$Ð$Ü—l‘l×0Ñ0°ÀÓKˆGò$÷$	ò 	ð Ð!8Ó8ð  $‘Ü˜E¤5×)Ñ)Ü>AÀ%ÈÔ>QÔRÒ>Q±d°a¸¡¨A¨qÖ 1Ñ>Q�ÑR�Ø 5Ó(ð
  $Ÿx™x×0Ñ0°·±¸uÏKÉKÐ0ÐX‘á*¨5°'Ó:�àÐ!5Ó5Ü! ,´×<Ñ<Ð<Ð<Ü$%§H¡H×$KÑ$KÓ$MÐ!Þ Ø*+¬C°Ó,>Ð+?¸uÀT×EVÑEVÐWhÓEiÐDjÐjkÐ(lÑ%ä(+ØŸ™×)Ñ)Ø ŸL™LØ.Ð/EÐ.FÀgÈlÈ^Ð[bÐcØ"3Ø".×"4Ñ"4ð	 *ð ó)Ð%ð &+°eÑ<¸^ÑL�ÙØ—L‘L *¨lÐ<Môð $5�
Ö ØÐ#3Ó3Ø×-×-à!%×!4Ñ!4Ø" N°E¸:ÀxÐQVó"’Jð "&×!=Ñ!=¸eÀUÓ!K’JØÐ#4Ó4Ü! ,´×9Ñ9Ð9Ð9Ø%1Ñ"��r˜6Ü"õ #à(,¯©ØŸ™ d¨B°¸¸Uô)ó#ó ’
ð  Ð#:Ó:ð "×BÑBÀ5È%ÓP’
ä! ,´×<Ñ<Ð<Ð<Ù*9Ø—L‘L ,°×0BÑ0Bó+Ñ'�˜ ð "ŸX™X×.Ñ.Ø—L‘L '°¸vð /ð ’
ð Ÿ(™(×+Ñ+Ø�J�<Ð Ø$Ü˜D×0Ñ0Ó2Ó3ð ,ð ˆKô
 —l‘l×6Ñ6°~ÀyÓQˆGØ×/Ñ/´¸wÓGˆGÜ˜g¤u×-Ñ-Ø! UÓ*Ø"&×"6Ñ"6Ó"8�KÜ˜{Ó+¨qÓ0Ð0Ð0á)4ó%Ú)4 ¸¸t»ÀsÈdÁ{Ÿ©ð "ð %ô ).¨mÓ(<�KÔ%Ø'(¨¯©°=Ó)AÐ(BÀ!Ð%D�NØ—I‘I×'Ñ'Ø&˜- {°>Ð2BÀ"ÀWÀIÈRÐPXÈzÐYZÐ[õð —I‘I×'Ñ'Ø&˜- {°4×3FÑ3FÓ3HÐ2IÈÈGÈ9ÐTVÐW_ÐV`Ð`aÐbôð Ð!5Ô5Ø&'¨
 |°6Ð$:Ð!Ø"Ÿm™m×>Ñ>Ó@�Ø—	‘	×#Ñ#Ø(Ð)¨°T×5HÑ5HÓ5JÐ4KÈ2Ü—{’{ ;Ó/×3Ñ3Ð4°B°t×7HÑ7HÈÓ7UÐ6VÐVWðYôð &+°eÑ<¸^ÑL�ð %Ñ0ð œ˜MÓ*Ð+¨5°×1BÑ1BÀ;Ó1OÐ0PÐPQÑRà2Ð3°5Ð9ð ð
 —‘×#Ñ#ðØ�˜WÐ%6Ð$7Ð7NÈwÈið XØ �M Ð$5Ð#6°b¸¸¸rÀ)Àð Mà�˜S¡¨{¨m¸5Ð,AÀ;Ó!OÐ Pð QØ"Ð# 3¡zÐ5FÐ4GÀuÐ2MÐO`Ó'aÐ&bð cðôñ  Ø×*Ñ*¨J¸ÐEVöô & n×5Ñ5Ø!×0Ñ0Ø °°zÀ8ÈUó’
ð  Ð#:Ô:Ø$% j \°Ð"6�Ø$% j \°Ð"6�ð —	‘	×#Ñ#Ø&Ð' {°4×3FÑ3FÓ3HÐ2IÐIZÐ[cÐZdÐdeÐfôð —	‘	×#Ñ#Ø&Ð' |°D×4GÑ4GÓ4IÐ3JÈ"ÈXÈJÐVWÐXôð —‘×#Ñ#ðØ$Ð% W¨_Ð,=ð >Ø(Ð)¨¨OÐ+<¸B¸u¸gÀRÌ×H\ÑH\ÐG]ð ^ðôð —‘×#Ñ#ðØ$Ð% S©°Ð6GÀuÐ4MÈÓ)_Ð(`ð aØ$Ð% S©°Ð6GÀuÐ4MÈÓ)_Ð(`ð aðôð (�
Ø!ŸX™XŸ_™_°5À
×@PÑ@P˜_ÐQ�
à!×GÑGØ×*Ñ*ØØØ#Ø#ØØó‘
ô  ×8Ò8¸ÈÓS�
Ù$ [°%Ó8�Ø! UÓ*Ø—L‘L×*Ñ*¨k¨]¸#¸g¸YÐ+GÕHà—L‘L×*Ñ*Ø&˜- s©:°g¸{Ó+KÐ*LÐMôð ¤§
¡
Ó*ð #'§(¡(×"3Ñ"3Ø×.Ñ.Ø&˜- |Ð4Ü#Ÿj™jØ)×/Ñ/ð	 #4ð #�Kñ 'Ø×*Ñ*¨J¸ÀTôð ×%×%Ð%Ü—l‘l×6Ñ6°~ÀyÓQˆGÜ#×-Ò-Ó/ˆJØ×.Ñ.°×0DÑ0DÓE�à—‘Ð.Ô/Ø×(Ñ(¨¯©«Ö6ñ Fð
 Ð!5Ó5Ø×&Ñ&×0Ñ0Ø!�l (¨4×+@Ñ+@ÀJÀ<ÈtÐATÓ+UÐ*VÐWôð  ×JÑJØ!�l %Ð(¨)°Wó�ð #Ÿm™m×>Ñ>Ó@�Ø×JÑJØ ¬U¯[ª[¸Ó-E×-IÑ-Ió�ñ   × 4Ñ 4°jÀ(ÈHÖUÜ% n×5Ñ5Ø%Ð)9Ó9Ð9Ð9Ø8BÑ5�˜Y¨Ø ×KÑKØÜ$ YÓ/Ø˜A‘Jó�	ð
 ×IÑIØÜ$ YÓ/Ø˜A‘Jó�ð
 #×MÑMØ!Ü$ YÓ/Ø˜A‘Jó�ð
 ×3Ñ3Ø×(Ñ(ØØØ!ØØØØØõ
ð  Ð#:Ó:Ø)3Ñ&�
˜JÜ! '¬8×4Ñ4Ð4Ð4Ø×JÑJØÔ 0°Ó ;¸WÀQ¹Zó�ð  ×JÑJØÔ 0°Ó ;¸WÀQ¹Zó�ð ×:Ñ:Ø×(Ñ(ØØØØØØõð ×GÑGØÔ 0°Ó ;¸Wó�ñ ' t×';Ñ';¸ZÈÐPTÔUØ×ÑÔà.8ˆ�‰× Ñ  Ñ+ä�j¤%×(Ñ(ÜÑLÁÓL×LÑLÐLÐLØ×"Ñ"×)Ñ)¨*Ô5ð Ð!LÓLÜ˜?Ó+¨qÓ0Ð0Ð0Ü"% j£/°OÑ"C�ä�z“?¤c¨/Ó&:Ó:Ð:Ð:Ü#& z°?Ö#C‘��ZØ!Ñ-Ð-Ð-Ø—9‘9 
Õ*Ø×*Ñ*×4Ñ4Ø˜%˜s 3 % tÔ,?À
Ó,KÐ+LÈAÐNöñ $Dð" Ðô ˜jÔ*;×<Ñ<Ð<Ð<Ø×"Ñ"×&Ñ& zÔ2ð ×Ñ ?°1Ñ#5Ó5Ø& qÑ)Ñ5Ð5Ð5Ø×&Ñ&×0Ñ0Ø!�l # j \°Ô6IÈ/ÐZ[ÑJ\Ó6]Ð5^Ð^_Ð`ôð Ðùòa Kùò4 Yùód  SùòR%s#   ªzÆzÆzÎzÙ5z Ú	z c                ó’  • U R                  XU5      nU R                  XU5      n[        S5       Vs/ s H%  n[        U R                  R	                  US95      PM'     snu  pxUR                  SU SU SU SU SU S[        R                   SU SU R                  U 5       SU SU R                  U 5       S35        Xx4$ s  snf )Nr!   r‚  ú
            r   ú9 = triton_helpers.online_softmax_reduce(
                ú)
            r  )	rÄ  rc  rÔ   rÆ   rg  rz  r"   r  r¼  )	r  r&  r$  r%  rg  rƒ  rW  r&  r'  s	            r†   Ú_online_softmax_reduceÚ#TritonKernel._online_softmax_reduce  sõ   € ð ×6Ñ6°vÐPUÓVˆØ×6Ñ6°vÐPUÓVˆÜMRÐSTÌXÓ!VÊXÈ¤# d§h¡h§o¡o¸E oÐ&BÖ"CÉXÑ!VÑˆ
Ø�‰ðØˆL˜˜:˜,ð 'Ø Ð!  OÐ#4°B°s°e¸2¼f×>RÑ>RÐ=Sð TØˆL˜˜D×1Ñ1°Z°LÓBÐCð DØˆL˜˜D×1Ñ1°Z°LÓBÐCð Dð	ô	
ð Ð%Ð%ùò "Ws   ²,Cc           
     óŽ  ^ ^^^• UUU 4S jX#U4 5       u  p#nSU SU SU ST S3	nU4S jnX#U4 V	s/ s H-  n	T R                   R                  TU" U	R                  5      S9PM/     n
n	TR                  SR	                  U
 Vs/ s H  n[        U5      PM     sn5       SU 35        [        U 4S jU
 5       5      $ s  sn	f s  snf )	z+
Helper to codegen triton_helpers.welford.
c              3  óJ   >#   • U  H  nTR                  TUT5      v •  M     g 7frš   )rÄ  )rö   rª  r&  rƒ  r  s     €€€r†   r  Ú(TritonKernel._welford.<locals>.<genexpr>   s*   øé € ð 
â+�ð ×(Ñ(¨°¸×>Ð>Ú+ùs   ƒ #ztriton_helpers.welford(r   r¡  c                ó2   >• [        U ST U TS-   S  -   5      $ )Nr   rN   )r×   )rÉ   rg  s    €r†   Úreduced_shapeÚ,TritonKernel._welford.<locals>.reduced_shape&  s$   ø€ Ü˜˜q ˜¨¨c°A©g¨iÐ(8Ñ8Ó9Ð9rˆ   r0  r  c              3  óZ   >#   • U  H   nTR                  XR                  5      v •  M"     g 7frš   )r¿  rÉ   )rö   rª  r  s     €r†   r  rA  /  s*   øé € ð 
â(�ð ×+Ñ+¨E·;±;×?Ð?Ú(ùó   ƒ(+)rÆ   rg  rÉ   rh  r   rÔ   r×   )r  r&  r  r  r  rg  rƒ  ÚwelfordrC  rª  Úwelford_resultsr  s   ``   ``     r†   r
  ÚTritonKernel._welford  sá   û€ ö
à FÑ+ó
Ñˆ�&ð ,¨D¨6°°B°4°r¸&¸ÀÀCÀ5ÈÐJˆõ	:ð
  FÑ+ó
â+�ð �H‰H�O‰O %©}¸U¿[¹[Ó/IˆOÓJÙ+ð 	ð 
ð 	×Ñ˜DŸI™I±Ó&G²°!¤s¨1¦v±Ñ&GÓHÐIÈÈWÈIÐVÔWäô 
á(ó
ó 
ð 	
ùò
ùò 'Hs   ³4B=ÂCc                ó˜  • U R                  5       U R                  -
  n[        U S3[        U R	                  5       5      U[
        R                  " 5       S9n[        U S3[        U R	                  5       5      U[
        R                  " 5       S9n	[        U S3[        U R	                  5       5      U[
        R                  " 5       S9n
U R                  R                  U SU R                  5        SU S35        U R                  R                  U	 SU R                  5        SU S35        U R                  R                  U
 SU R                  5        SU S35        US:X  a=  Uu  p¼nU R                  R                  S	U S
U	 S
U
 SU SU	 SU
 SU SU SU S35        O9US:X  d   eU R                  R                  S	U S
U	 S
U
 SU SU SU	 SU
 S35        U R                  R                  SU SU" U S3U5       SU	 SU" U	 S3U	5       SU
 SU" U
 S3U
5       S35        UnU R                  U R                  USSUU	U
UU5	      $ )z%Helper to codegen a welford reductionÚ_mean©rÉ   rƒ  r›  Ú_m2Ú_weightr   r   r¡  r÷  rà  rû  z<_next = triton_helpers.welford_combine(
                    z,
                    rü  rö  z;_next = triton_helpers.welford_reduce(
                    z1, roffset == 0
                )
                z            r  rý  r:  N)rÑ   rX  r–  r×   rF  r    Úunknownre  rh  rÞ  r5  rz  r  rA  )r  rQ  rƒ   rª  r9  r  rƒ  rg  r#  Úaccumulator_m2Úaccumulator_weightr  r  r  r-  s                  r†   rö  ÚTritonKernel.welford_reduce4  sþ  € ð ×%Ñ%Ó'¨$×*AÑ*AÑAˆä'Øˆl˜%Ð Ü˜×,Ñ,Ó.Ó/ØÜ×&Ò&Ó(ñ	
ˆô +Øˆl˜#ÐÜ˜×,Ñ,Ó.Ó/ØÜ×&Ò&Ó(ñ	
ˆô /Øˆl˜'Ð"Ü˜×,Ñ,Ó.Ó/ØÜ×&Ò&Ó(ñ	
Ðð 	�	‰	×ÑØˆm˜<¨×(;Ñ(;Ó(=Ð'>¸bÀÀ
È!ÐLô	
ð 	�	‰	×ÑØÐ˜l¨4×+>Ñ+>Ó+@Ð*AÀÀHÀ:ÈQÐOô	
ð 	�	‰	×ÑØ!Ð" ,¨t×/BÑ/BÓ/DÐ.EÀRÈÀzÐQRÐSô	
ð Ð.Ó.Ø$ÑˆD�fØ�L‰L×ÑðØ�˜W ^Ð$4°GÐ<NÐ;Oð PØ �M  NÐ#3°2Ð6HÐ5Ið JØ�F˜"˜R˜D  6 (ð +ðõð "Ð%5Ó5Ð5Ð5Ø�L‰L×ÑðØ�˜W ^Ð$4°GÐ<NÐ;Oð PØ�G˜2˜k˜]¨"¨^Ð,<¸BÐ?QÐ>Rð Sðôð 	�‰×ÑðØˆM˜™Z¨;¨-°uÐ(=¸{ÓKÐLð MØÐ˜C¡
¨nÐ-=¸UÐ+CÀ^Ó TÐUð VØÐ  ¡JÐ2DÐ1EÀUÐ/KÐM_Ó$`Ð#að bðô	
ð !ˆØ×2Ñ2Ø×"Ñ"ØØØØØØØØó

ð 
	
rˆ   c
           
     ó
  • [        U R                  XXgX‰5      5      n
X#U/n[        [        Xº5      5       HA  u  nu  nu  pïUc  U R                  R                  XŸS9nXÛU'   UR                  U SU 35        MC     [        U5      $ )z0Helper to codegen call to triton_helpers.welfordr0  r  )ra  r
  r–  r¤  rÆ   rg  rz  r×   )r  r&  r-  r.  r/  r  r  r  rg  rƒ  r¿  Úresult_exprsrU  Úresult_exprrª  rÉ   s                   r†   r  Ú+TritonKernel.welford_reduce_final_reduction|  s‹   € ô �d—m‘m F°"¸cÓIÓJˆà#°Ð>ˆÜ09¼#¸lÓ:SÖ0TÑ,ˆAÑ,�™^˜eØÑ"Ø"Ÿh™hŸo™o°E˜oÐG�Ø"-˜Q‘Ø�M‰M˜[˜M¨¨U¨GÐ4Ö5ñ	 1Uô �\Ó"Ð"rˆ   c                ó  • U R                  XU5      nU R                  XU5      n	UR                  SU SU SU SU	 SU S[        R                   SU SU R	                  U 5       SU SU R	                  U 5       S35        X#4$ )Nr:  r   r;  r<  r  )rÄ  rz  r"   r  r¼  )
r  r&  r&  r'  r3  r4  rg  rƒ  r$  r%  s
             r†   r  Ú2TritonKernel.online_softmax_reduce_final_reduction”  s½   € ð ×6Ñ6°vÈÓOˆØ×6Ñ6°vÈÓOˆØ�‰ðØˆL˜˜:˜,ð 'Ø Ð!  OÐ#4°B°s°e¸2¼f×>RÑ>RÐ=Sð TØˆL˜˜D×1Ñ1°Z°LÓBÐCð DØˆL˜˜D×1Ñ1°Z°LÓBÐCð Dð	ô	
ð Ð%Ð%rˆ   c                óN   • U R                   (       a  U R                   S   $ [        $ )NÚRSPLIT)r  r3   r  s    r†   Ú
max_rsplitÚTritonKernel.max_rsplit£  s"   € Ø××Ø×$Ñ$ XÑ.Ð.Ü Ð rˆ   c                ó  • U R                   S   nU R                  5       (       d  SOSnXBR                  -  U R                  5       -  nU R                  R                  U5      u  pxU R                  R                  SU SU SU R                  U5       S[        U5       SU S	U S
U S3SS9  U R                  U S3SS/U[        R                  " 5       S9n	U R                  R                  U	 SU S[        U5       S35        U	$ )zé
Generate code to save a [XBLOCK, RSPLIT] temporary workspace, where each thread block writes a different
column.  After the barrier, every thread block loads the completed value so that it can compute the final
value independently.
rN  zxindex < xnumelNrã  z_ws = (r`  z).to(tl.pointer_type(z))
                tl.store(z%_ws + (xindex * RSPLIT + rsplit_id), r   r<  T©ÚstripÚ_peersrñ   rZ  rL  z = tl.load(z_ws + (xindex * RSPLIT + rsplit_arange), rsplit_mask, eviction_policy='evict_first', other=triton_helpers.if_mask(rsplit_mask, r@  )rZ  r~  r‹  r[  ry  ræ  rA  rz  r§  rF   Úcreate_cse_varr    rO  rB  rh  r`   )
r  rQ  rƒ  Údefault_valÚxnumelr  râ  rä  rå  r5  s
             r†   r  Ú7TritonKernel.codegen_cooperative_reduction_peer_combine¨  sR  € ð —‘˜SÑ!ˆØ(,×(@Ñ(@×(BÑ(BÑ ÈˆàŸ.™.Ñ(¨4¯?©?Ó+<Ñ<ˆØ!×GÑG×PÑPÐQWÓXÑˆà×Ñ×%Ñ%ðØ�˜G G 9¨C°×0AÑ0AÀ)Ó0LÐ/MÐMbÔcnÐotÓcuÐbvð wØ$˜Ð%JÈ:È,ÐVXÐY]ÐX^ð _ðð ð 	&ñ 	
ð ×#Ñ#Øˆl˜&Ð!Ø˜XÐ&ØÜ×&Ò&Ó(ð	 $ð 
ˆð 	×Ñ×&Ñ&Øˆg�[  ð -eÜerÐs~Óeð  eAð  ACðDô	
ð ˆrˆ   c                ó|  • U R                   (       d   eSU l         [        R                  R                  U5      nU R	                  USU R                  U USSS9S9nSU l         U R                  R                  U5      n[        R                  " 5       nU R                  (       a*  UR                  U R                  XR                  5      5        U R                  U R                  U5        [        U[         ["        45      (       aZ  U R                  R%                  ['        UU R)                  UUUR+                  U5      USUR-                  5       < 35      5      5        OÏ[        U[.        5      (       d   eUR0                  n[3        U5      (       a`  UR4                  bS  [7        S UR4                   5       5      (       d2  SR9                  [;        [<        UR4                  5      5      n	USU	 S	3-  nU R                  R%                  ['        US
U SU SU SUR>                   S	3	5      5        URA                  5         g )NFT)rÅ   rƒ  r­  r  r@  r  c              3  ó>   #   • U  H  n[        U5      S :H  v •  M     g7fr‚  rÜ  r:  s     r†   r  Ú/TritonKernel.store_reduction.<locals>.<genexpr>û  r‡  r„  r   r€  r¡  r  r½  r	  )!rY  rL   rP  r  rD  rm  rÂ  rO  r‹  rŒ  rQ  r�  rŽ  rB  rK  rŠ   rÈ  r»  rh  rU   r  rÃ  r‰  rø   rù   r“   rÉ   rR  r   r"  rÔ   r#  r�  )
r  rÈ   rþ   rª  rƒ  rD  rÚ   r”  r’  r“  s
             r†   Ústore_reductionÚTritonKernel.store_reductionÉ  s	  € ð ×$×$Ð$Ð$Ø %ˆÔÜ—‘×!Ñ! $Ó'ˆØ—=‘=ØØØ&*×&HÑ&HØØØØð	 'Ið 'ð !ð 	
ˆð !%ˆÔØ�i‰i×Ñ˜tÓ$ˆä×)Ò)Ó+ˆ
Ø×%×%Ø×$Ñ$Ø×,Ñ,¨T×3GÑ3GÓHôð 	×&Ñ& t×';Ñ';¸SÔAä�h¤Ô2IÐ J×KÑKØ× Ñ ×*Ñ*ÜØØ×5Ñ5ØØ Ø Ÿ™¨Ó,ØØ+¨H×,CÑ,CÓ,EÑ+HÐIóó	õô ˜h¬×8Ñ8Ð8Ð8à#×-Ñ-ˆLä% e×,Ñ,Ø—K‘KÑ+ÜÑ?°5·;²;Ó?×?Ñ?à"Ÿi™i¬¬C°·±Ó(=Ó>�Ø .°°¸QÐ ?Ñ?�à× Ñ ×*Ñ*ÜØØ ˜u D¨¨°c¸%¸ÀÀ8×CTÑCTÐBUÐUVÐWóôð 	×ÑÕrˆ   c                óf  ^^^^^^^• [        5       mTR                  S5        [        5       m[        S5       V^s/ s H+  m[	        UU4S j[        [        X#5      5       5       5      PM-     nnSR                  S [        R                  R                  U5       5       5      nTR                  SU S35        [        5       mSmS	S
KJn  S	SKJn  U" 5       mU" 5       m " UUUUUU4S jS[         5      n	TR#                  5          [$        R&                  " U	" 5       5         U" U6 n
SR                  S U
 5       5      n
TR                  SU
 35        S S S 5        S S S 5        U R(                  R+                  TR-                  5       TS9$ s  snf ! , (       d  f       NC= f! , (       d  f       NL= f)Nz@triton.jitr!   c              3  ón   >#   • U  H*  u  nu  p#TR                  S T SU 3X2R                  S9v •  M,     g7f)r¡  rW  r0  N)r  rÉ   )rö   r¡  rª  rƒ  rÆ   rU  s       €€r†   r  Ú,TritonKernel._lift_helper.<locals>.<genexpr>  s<   øé € ð â)GÑ%�A‘~˜ð —‘˜s 1 # Q q c˜]°%¿{¹{�ÕKÚ)Gùs   ƒ25r   c              3  ó8   #   • U  H  n[        U5      v •  M     g 7frš   rÜ  r:  s     r†   r  rl    s   é € ÐRÒ.Q¨œc !Ÿf˜fÒ.Qùrh  zdef {name}(ú):r~  r   rn   )ÚShapePropagationOpsHandlerc                  ó@   >• \ rS rSr        SU UUUUU4S jjrSrg)Ú+TritonKernel._lift_helper.<locals>.CSEProxyi*  c                ó    >• T	SU 3-  m	[        TU5      " U0 UD6n[        TU5      " U0 UD6nTR                  T[        T
U5      " U0 UD6UUS9$ )NrW  r0  )rÀ  r4  )r  rÈ   rÂ  r   Úoutput_dtypeÚoutput_shaperÆ   Údtype_handlerÚhelperÚhelper_nameÚ	overridesÚshape_handlers         €€€€€€r†   Ú_defaultÚ4TritonKernel._lift_helper.<locals>.CSEProxy._default+  s”   ø€ ð   4 &˜zÑ)�ä&Ø!Øô ð ð #ð "ñ #�ô
  'Ø!Øô ð ð #ð "ñ #�ð
 —|‘|ØÜ˜I tÔ,¨dÐ=°fÑ=Ø&Ø&ð	 $ð ð rˆ   r‚   N)rÈ   rÔ   rÂ  ztuple[Any, ...]r   rt  r¦   r	   )r›   r¨   r©   rª   rz  r®   )rÆ   ru  rv  rw  rx  ry  s   €€€€€€r†   ÚCSEProxyrq  *  s-   ø† ðØðØ'6ðØ@Nðà÷÷ rˆ   r|  c              3  ó8   #   • U  H  n[        U5      v •  M     g 7frš   rÜ  )rö   rO  s     r†   r  rl  D  s   é € ÐBº'°¤ F§ º'ùrh  úreturn r  )rV   rh  rS   rc  r×   r–  r¤  r   r½  r¾  Úfrom_iterablerÐ  r¨  ro   Ú!torch._inductor.shape_propagationro  r,   rž  rL   Úset_ops_handlerr5  r„  Úgetvalue)r  r;  r¿  ÚdtypesrU  rÂ  Ú	signaturero   ro  r|  ÚoutputsrÆ   ru  rv  rw  rx  ry  s       `      @@@@@@r†   Ú_lift_helperÚTritonKernel._lift_helper	  sg  þ€ ô
  Ó!ˆØ×Ñ˜Ô'Ü‹eˆô ˜1”Xô
ò
 �ô	 õ ä)2´3°vÓ3FÔ)Góö ñ ð 	ð 
ð —I‘IÑR¬i¯o©o×.KÑ.KÈDÔ.QÓRÓRˆ	Ø×Ñ˜=¨¨°2Ð6Ô7ä#Ó%ˆ	ð *ˆåPÝPá2Ó4ˆÙ2Ó4ˆ÷	ô 	”~ô 	ð0 �]‰]�_œa×/Ò/±³
Õ;Ù˜$�iˆGØ—i‘iÑB¹'ÓBÓBˆGØ×Ñ˜w w iÐ0Ô1÷ <�_ð
 ×$Ñ$×(Ñ(¨¯©Ó):ÀkÐ(ÐRÐRùòk
÷` <Õ;ú�_�_ús)   »2FÄF"Ä"2FÅF"Æ
F	ÆF"Æ"
F0c                ó€  ^ ^• T R                   (       d   eT R                  (       a   S5       e[        S T R                   5       5      nT R	                  U5        [        U5      nT R                  (       a   S5       e/ n/ n[        S U 5       5      n[        R                  " T R                  R                  T R                  5      nT R                  X#U5      nT R                  5       T R                  -
  n	[!        X15       GHV  u  p«T R                  R                  T R                  U
 S[#        U5       S3UU
R$                  S9nT R                  R                  T R                  SU S	T R'                  5        S3U[        T R)                  5       5      S9n
UR+                  U
5        [-        U5      nT R.                  (       a  MÆ  T R)                  5       nS
US'   T R                  R1                  X¾S9nSS	R3                  U5       S3nUR4                  (       a  SOSnT R6                  R9                  U SU S	U S	U S35        UR+                  U5        GMY     S mUU 4S jnU" ST" U5       SU	 S	U S3UUU5      nT R.                  (       dä  S nU Vs/ s H(  nU" SU S3[;        UR<                  5      U" U5      S9PM*     nnU" [        U5      [        U5      5      nU" [        U5      U5      n[!        UU5       VVs/ s H)  u  nnU" SU S	U S3UR<                  UR$                  S9PM+     nnn[!        UUU5       H+  u  nnnT R                  R9                  U SU S	U S35        M-     OUnU H*  n[?        U[@        5      (       d   e[        U5      Ul!        M,     [        U5      $ s  snf s  snnf )z*
Perform an associative scan on 'values'.
ÚTODOc              3  ó>   #   • U  H  oR                    S 3v •  M     g7frâ  rÍ  r„  s     r†   r  Ú$TritonKernel.scan.<locals>.<genexpr>V  rÎ  r„  z(ops.scan not supported inside ops.maskedc              3  ó8   #   • U  H  n[        U5      v •  M     g 7frš   ©rH   ©rö   rƒ  s     r†   r  r‹  ^  ó   é € ÐFºv°eÔ*¨5×1Ð1ºvùrh  r»  r¡  r0  rŸ  r   rÃ   rK  rÖ  r×  zfloat('nan')z-1rù  c                ó2   • SR                  S U  5       5      $ )Nr?  c              3  ó*   #   • U  H	  o S 3v •  M     g7f©rè  Nr‚   ©rö   rª  s     r†   r  Ú1TritonKernel.scan.<locals>.csv.<locals>.<genexpr>‚  ó   é € Ð<²V¨E˜g Q�K²Vùr\  ©r   ©r¿  s    r†   ÚcsvÚTritonKernel.scan.<locals>.csv�  ó   € Ø—8‘8Ñ<±VÓ<Ó<Ð<rˆ   c           	     ód  >• [        U5      n[        U5       Vs/ s H  oP SU SU 3PM     nn[        U4S jU 5       5      (       a,  U Vs/ s H  nTR                  R	                  U5      PM      sn$ [        X15       VV	s/ s H(  u  p‰TR                  R                  X‰R                  S9PM*     n
nn	TR                  R                  T" U
5       SU  35        [        X¦5       H-  u  p·U(       a  X+l
        TR                  R                  X{5        M/     [        U
5      $ s  snf s  snf s  sn	nf )Nr   c              3  óZ   >#   • U  H   nTR                   R                  U5      v •  M"     g 7frš   ©rÆ   Úcontains©rö   rj  r  s     €r†   r  Ú:TritonKernel.scan.<locals>.cse_multiple.<locals>.<genexpr>‡  ó#   øé € ÐLÂ°I�4—8‘8×$Ñ$ Y×/Ð/ÂùrF  r0  r  ©r�   rc  rR  rÆ   rè  r¤  rg  rÉ   r5  rh  rû   ri  r×   )r%  r¿  r¢  rƒ  r¡  rU  Ú
cache_keysrj  rƒ  rª  Úresult_varsrQ  r˜  r  s               €€r†   Úcse_multipleÚ'TritonKernel.scan.<locals>.cse_multiple„  s  ø€ Ü�F“ˆAÜ;@À¼8ÓDº8°a˜F " Q C r¨%¨Ó1¹8ˆJÐDÜÔLÁÓL×LÑLÙAKÓLÂ°I˜Ÿ™Ÿ™ YÖ/ÁÑLÐLô '*¨&Ô&9ôâ&9‘N�Uð —‘—‘ e·;±;�Ó?Ù&9ð ñ ð �L‰L×"Ñ"Ù�{Ó#Ð$ C¨ vÐ.ôô *-¨[Ö)EÑ%�
ÞØ+0Ô(Ø—‘—‘˜YÖ3ñ *Fô ˜Ó%Ð%ùò EùâLùóó   šD"Á%D'Â/D,ztl.associative_scan((r	  c                óV   • U R                   c  g [        U R                   5      nSUS'   U$ )NrÃ   rK  )rÉ   ra  )rÚ   rÉ   s     r†   Ú_partial_scan_shapeÚ.TritonKernel.scan.<locals>._partial_scan_shape¡  s*   € Ø—9‘9Ñ$Øä  §¡›O�EØ #�E˜"‘IØ �Lrˆ   ztriton_helpers.select_one((z1), rbase == (RBLOCK - 1), dim=-1, keep_dims=True)ztl.where(roffset > 0, z = tl.where(roffset > 0, )"rY  rQ  r   re  rÎ  r!  rü  r×   r6  r7  rÆ   r4  r5  r†  rÑ   rX  r¤  r„  rÉ   rÞ  rF  rØ  r�  r  rg  r   r‘  re  rh  rH   rƒ  rŠ   r–  rû   ) r  rƒ  r(  r¿  r¢  Úbroadcasted_valuesÚaccumulatorsÚcse_computeÚcombine_helper_fnrg  rª  rƒ  Úvalue_dtyper  Úreduced_sizer#  Úreduced_size_strr6  r¥  Úpartial_scan_varsr©  Úpartial_scan_varÚpartial_reduce_varsÚ	accs_nextÚfull_scan_varsÚ	full_scanÚpartial_scanr¤  Úacc_nextÚpartial_reducerQ  r˜  s    `                              @r†   ÚscanÚTritonKernel.scanI  sò  ù€ ð ×$×$Ð$Ð$Ø×-×-Ð5¨vÓ5Ð-ÜÑM¸D×<LÒ<LÓMÓMˆØ×Ñ˜%Ô Ü�u“ˆØ—?—?ÐNÐ$NÓNÐ"àÐØˆäÑF¹vÓFÓFˆÜ×'Ò'¨¯©×(9Ñ(9¸4¿<¹<ÓHˆØ ×-Ñ-¨jÀ&ÓIÐØ×%Ñ%Ó'¨$×*AÑ*AÑAˆä ×/‰LˆEØŸ(™(×+Ñ+Ø—‘Ø�'˜Ô1°%Ó8Ð9¸Ð;ØØ—k‘kð	 ,ð ˆKð —H‘H×%Ñ%Ø—‘Ø" ; -¨r°$×2EÑ2EÓ2GÐ1HÈÐJØÜ˜D×0Ñ0Ó2Ó3ð	 &ð ˆEð ×%Ñ% eÔ,ä& uÓ-ˆHà×,×,Ñ,Ø#×3Ñ3Ó5�Ø#&�˜RÑ Ø"Ÿh™hŸo™o°E˜oÐN�Ø%& t§y¡y°Ó'>Ð&?¸qÐ#AÐ à,1×,C×,C™.È�Ø—	‘	×#Ñ#Ø"�m ;Ð/?Ð.@ÀÀ7À)È2ÈhÈZÐWXÐYôð ×#Ñ# K×0ñ9 0ò<	=ö	&ñ$ )Ø#¡CÐ(:Ó$;Ð#<¸CÀ¸uÀBÐGXÐFYÐYZÐ[ØØØó	
Ðð ×(×(ò!ñ ):ó#ò ):Ð$ñ Ø1Ð2BÐ1CÐCtÐuÜ-Ð.>×.DÑ.DÓEÙ-Ð.>Ó?ôñ
 ):ð  ð #ñ #¤5¨Ó#6¼Ð>QÓ8RÓSˆIÙ'¬¨lÓ(;Ð=NÓOˆNô 03°>ÐCTÔ/Uôò 0VÑ+�I˜|ñ Ø,¨Y¨K°r¸,¸ÀqÐIØ&×,Ñ,Ø&×,Ñ,ôñ
 0Vð ñ ô :=Ø˜<Ð)<ö:Ñ5�˜+ ~ð —‘×&Ñ&Ø"�mÐ#<¸X¸JÀbÈÐHXÐXYÐZöò:ð ,ˆKã%ˆJÜ˜jÔ*;×<Ñ<Ð<Ð<Ü#-¨eÓ#4ˆJÖ ñ &ô �[Ó!Ð!ùò?#ùós   Ê/N5Ì0N:c                óh  ^ ^• T R                   (       d   eT R                  (       a   S5       e[        S T R                   5       5      nT R	                  U5        [        U5      nT R                  (       a   S5       eT R                  (       d   S5       e[        R                  " T R                  R                  T R                  5      nT R                  5       T R                  -
  n[        S U 5       5      n[!        U5      [!        U5      :X  d   e[#        U5       VV	s/ s H;  u  p‰U" SU	 ST R%                  5        S3X   [        T R'                  5       5      S	9PM=     n
nn	S
 mUU 4S jnT R                  S   R(                  (       d   eT R+                  T R                  S   5      (       a  SOSn[!        U5      S:X  a%  SU
S    SU
S    SU SU SU SU S3nU" XÚXQ5      nO[-        S5      e[/        Xâ5       H  u  nnX_l        UR2                  Ul        M     [        U5      $ s  sn	nf )Nr‰  c              3  ó>   #   • U  H  oR                    S 3v •  M     g7frâ  rÍ  r„  s     r†   r  Ú$TritonKernel.sort.<locals>.<genexpr>Ó  rÎ  r„  z(ops.sort not supported inside ops.maskedz3ops.sort is only supported in persistent reductionsc              3  ó8   #   • U  H  n[        U5      v •  M     g 7frš   r�  rŽ  s     r†   r  r¿  Þ  r�  rh  rŸ  r   r¡  r0  c                ó2   • SR                  S U  5       5      $ )Nr?  c              3  ó*   #   • U  H	  o S 3v •  M     g7fr’  r‚   r“  s     r†   r  Ú1TritonKernel.sort.<locals>.csv.<locals>.<genexpr>ê  r•  r\  r–  r—  s    r†   r˜  ÚTritonKernel.sort.<locals>.csvé  rš  rˆ   c           	     ód  >• [        U5      n[        U5       Vs/ s H  oP SU SU 3PM     nn[        U4S jU 5       5      (       a,  U Vs/ s H  nTR                  R	                  U5      PM      sn$ [        X15       VV	s/ s H(  u  p‰TR                  R                  X‰R                  S9PM*     n
nn	TR                  R                  T" U
5       SU  35        [        X¦5       H-  u  p·U(       a  X+l
        TR                  R                  X{5        M/     [        U
5      $ s  snf s  snf s  sn	nf )Nr   c              3  óZ   >#   • U  H   nTR                   R                  U5      v •  M"     g 7frš   r�  rŸ  s     €r†   r  Ú:TritonKernel.sort.<locals>.cse_multiple.<locals>.<genexpr>ï  r¡  rF  r0  r  r¢  )r%  r«  r¢  rƒ  r¡  rU  r£  rj  rƒ  rª  r¤  rQ  r˜  r  s               €€r†   r¥  Ú'TritonKernel.sort.<locals>.cse_multipleì  s  ø€ ÜÐ&Ó'ˆAÜ;@À¼8ÓDº8°a˜F " Q C r¨%¨Ó1¹8ˆJÐDÜÔLÁÓL×LÑLÙAKÓLÂ°I˜Ÿ™Ÿ™ YÖ/ÁÑLÐLô %(¨Ô$Côâ$C‘L�Eð —‘—‘ e·;±;�Ó?Ù$Cð ñ ð �L‰L×"Ñ"Ù�{Ó#Ð$ C¨ vÐ.ôô *-¨[Ö)EÑ%�
ÞØ+0Ô(Ø—‘—‘˜YÖ3ñ *Fô ˜Ó%Ð%ùò EùâLùór§  rK  r§   Úrnumelr!   ztriton_helpers.sort_with_index(r   rN   z	, stable=z, descending=zUnhandled sort)rY  rQ  r   re  rÎ  r!  rü  r  r6  r7  rÆ   r4  r5  rÑ   rX  r×   r�   r–  rÞ  rF  r\  r'  r`  r¤  rû   r›  )r  rƒ  r¿  ÚstableÚ
descendingr¢  r­  rg  rU  rª  r«  r¥  rÉ  r%  r¤  rQ  Ú	input_varr˜  s   `                @r†   r£  ÚTritonKernel.sortÊ  s5  ù€ ð ×$×$Ð$Ð$Ø×-×-Ð5¨vÓ5Ð-ÜÑM¸D×<LÒ<LÓMÓMˆØ×Ñ˜%Ô Ü�u“ˆØ—?—?ÐNÐ$NÓNÐ"Ø×(×(ð 	
ØAó	
Ð(ô  ×'Ò'¨¯©×(9Ñ(9¸4¿<¹<ÓHˆØ×%Ñ%Ó'¨$×*AÑ*AÑAˆäÑF¹vÓFÓFˆÜ�6‹{œc &›kÓ)Ð)Ð)ô & fÔ-ô
ò .‘�ñ Ø" 5 '¨¨D×,?Ñ,?Ó,AÐ+BÀ!ÐDØ‘iÜ˜D×0Ñ0Ó2Ó3ôñ
 .ð 	ñ 
ò	=ö	&ð$ ×Ñ Ñ#×0×0Ð0Ð0Ø×2Ñ2°4×3CÑ3CÀBÑ3G×HÑH‘Èhˆäˆv‹;˜!Óà1Ð2DÀQÑ2GÐ1HÈÐK]Ð^_ÑK`ÐJað bØ�8˜2˜c˜U )¨F¨8°=ÀÀÈAðOð ñ ' tÀÓO‰Kä Ð!1Ó2Ð2ä%(¨Ö%=Ñ!ˆJ˜	Ø#(Ô Ø )× 0Ñ 0ˆJÖñ &>ô �[Ó!Ð!ùó]
s   Ä"AH.c                óÆ   • U R                   (       d  gUR                  U R                   5        U R                   R                  5         U R                  R                  5         g)z„
Generate the output from prologue. This should be
extracted from the subgraph, which is why this is
partitioned from codegen_body.
N)r@  rz  Úclearr?  )r  rV  s     r†   Úcodegen_prologueÚTritonKernel.codegen_prologue  s@   € ð �}�}Øà�‰�D—M‘MÔ"Ø�‰×ÑÔØ×Ñ×!Ñ!Õ#rˆ   c                ó<  • U R                   (       dV  U R                  (       dE  U R                  (       d4  U R                  (       d#  U R                  (       d  U R
                  (       d  gU R                   Vs/ s H  oR                  (       d  M  UPM     nnU R                  (       GaÑ  U R                  (       d   S5       e0 n[        U R                  5       H¦  u  pEUR                  n[        R                  R                  U[         R"                  5      nU R%                  [&        U5      nSU 3nU R(                  R+                  U SU S35        U R,                  R/                  U[         R"                  SS9X8'   M¨     U R(                  R+                  S5        U R(                  R+                  S	5        U R(                  R1                  S
S9   U R3                  5       (       d[  U R                  S   n	U	R4                  S:X  d   eU	R4                  n
U R(                  R+                  U
 SU	R6                   SU
 S35        U R(                  R9                  U R                   5        U R(                  R;                  S/5        U R(                  R9                  U R                  5        U R(                  R9                  U R                  5        U R(                  R9                  U R                  5        U R(                  R9                  U R
                  5        [        U R                  5       Hù  u  pEUR<                  nSU 3n[        R>                  " UR                  [         R"                  5      n[A        UR                  5      nU R,                  RC                  U R(                  U SU S3URD                  SS9nSSK#nURH                  RJ                  RM                  U SU R(                  5         U" X8   U5      nSSS5        U R(                  R+                  U SW 35        Mû     SSS5        [O        [Q        U R                  5      5       H%  nU R(                  R+                  SU SU S35        M'     GO;U RR                  (       Ga•  [Q        U5      S:”  Ga…  [        U5       GH  u  nnU R(                  R1                  US9   UR4                  nU RT                  (       a  SOSnU RT                  (       a  SOU S3n[         RV                  RX                  (       a  [[        5       S:”  a  SnOSnU R(                  R+                  S U S!U S"U S"UR]                  5        S#U S$35        SSS5        U R(                  R1                  US
-   S9   U R_                  XR(                  5        SSS5        GM     U R(                  R1                  [Q        U5      S9   U Ra                  U R(                  5        U R(                  R9                  U R                   5        U R(                  R9                  U R                  5        U R(                  R9                  U R                  5        U R(                  R9                  U R                  5        SSS5        [c        / [        U5      Q5       GHr  u  nnU R(                  R1                  US
-   S9   U Rd                  URf                     Ri                  5        Hæ  u  nnU[Q        U5      S
-
  :  av  UUS
-      nU Rd                  URf                     U   n[j        Rm                  U5      n[o        URp                  U5      n[s        UU5       VVs/ s H  u  nnUUU-  -
  PM     nnnU R(                  R+                  [u        U Rv                  U   U S%U S"[x        Rz                  R}                  U5       S&35      5        Mè     SSS5        U R,                  R                  U R€                  5        URƒ                  5         GMu     O”U R(                  R9                  U R                   5        U R(                  R9                  U R                  5        U R(                  R9                  U R                  5        U R(                  R9                  U R                  5        U R(                  R9                  U R                  5        U RT                  (       ai  U R                  (       d  U R
                  (       aG  U R„                   S'3nU R(                  R9                  S(U S)3S*S+9  U R†                  R‰                  5         U R                  (       d%  U R(                  R9                  U R
                  5        U R                   R‹                  5         U R                  R‹                  5         U R                  R‹                  5         U R                  R‹                  5         U R                  R‹                  5         U R
                  R‹                  5         gs  snf ! , (       d  f       GN6= f! , (       d  f       GN= f! , (       d  f       GNë= f! , (       d  f       GMË  = f! , (       d  f       GNñ= fs  snnf ! , (       d  f       GN¶= f),zØ
Concat output code from index_code, loads, compute, stores,
suffix into self.body.

For pointwise kernels, this is called just once at the end.

For reduction kernels, this generates a loop over the reduction
axis.
Nz1Mix order reduction requires persistent reductionÚaccumz = tl.full([R0_BLOCK], z, tl.float32)[None, :])rÃ   ÚR0_BLOCKr0  z/split_size = min(RSPLIT_SIZE, xnumel - xoffset)z@for _ in tl.range(0, split_size, XBLOCK, num_stages=NUM_STAGES):rN   )rº   r   rN  úmask = ú < r  zxindex += XBLOCKr)  z, 0))rÔ  r5  r  z&tl.store(ws_ptr + (tl.program_id(0) + z3 * tl.num_programs(0)) * r0_numel + r0_index, accumz
, r0_mask)Úrsplit_startr  Ú
rsplit_end)r   r!   z, num_stages = 2r°   zfor zoffset in tl.range(r   r¾   rn  z = tl.advance(r¡  z + tl.program_id(1)zR
                if HAS_RSPLIT:
                    triton_helpers.x_grid_barrier(rá  Tr^  )FÚindexing_coder*  r‰  r5  rA  rB  re  rƒ  Úmix_order_reductionr  r–  r^  rƒ   r#   r1  r  rœ   r,  rñ  r`   re  rh  rÆ   r  rž  r~  rÓ   rÈ   rz  Ú
writelinesrª  r  r‡   r4  rƒ  ÚunittestÚmockÚpatchÚobjectrc  r�   rY  rQ  rð  rñ  r   rõ   r†  rŠ  rb  rH  rÝ   r¼  r¸   rÕ   r   r  r¤  rU   rG  rL   rÅ   r§  Ú
invalidaterC  Úcache_clearrx  ry  rê  rÏ  )r  rá   Ú
loop_treesÚaccumname2varr…  Úpartial_accumrƒ   r6  rÈ   ÚentryrN  rÚ   r(  Útriton_reduction_functionÚnewvalrÜ  r)  ÚlevelrÓ   Ú
loop_startÚloop_endÚ
num_stagesrA  ÚadvancementÚ	prev_treeÚprev_advancementÚ
prev_blockÚprev_num_iterÚcurÚprevÚsem_ptrs                                  r†   Úcodegen_bodyÚTritonKernel.codegen_body  s‡  € ð ××Ø�z�zØ�{�{Ø�|�|Ø×%×%Ø×#×#àà'+×'7Ò'7ÓHÒ'7˜t¿<½<—dÑ'7ˆ
ÐHØ×#×#Ð#Ø×,×,ð ØCóÐ,ð ˆMÜ&/°×0MÑ0MÖ&NÑ"�Ø!.×!=Ñ!=�ÜŸ,™,×:Ñ:¸>Ì5Ï;É;ÓW�Ø×3Ñ3´MÀ7ÓK�Ø˜s˜e�}�Ø—	‘	×#Ñ#Ø�fÐ3°G°9Ð<RÐSôð '+§h¡h×&7Ñ&7Ø¤§¡Ð3Dð '8ð '�Ó#ñ 'Oð �I‰I×ÑÐ QÔRØ�I‰I×ÑØRôð —‘×!Ñ!¨Ð!Ò+à×/Ñ/×1Ñ1Ø ×,Ñ,¨QÑ/�EØ Ÿ<™<¨3Ó.Ð.Ð.ØŸ™�AØ—I‘I×'Ñ'¨1¨#¨W°U·Z±Z°LÀÀAÀ3ÀeÐ(LÔMØ—	‘	× Ñ  ×!3Ñ!3Ô4Ø—	‘	×$Ñ$à*ðôð
 —	‘	× Ñ  §¡Ô,Ø—	‘	× Ñ  §¡Ô.Ø—	‘	× Ñ  §¡Ô-Ø—	‘	× Ñ  ×!5Ñ!5Ô6ô +4°D×4QÑ4QÖ*RÑ&�CØ'×-Ñ-�CØ" 3 %˜=�DÜ!#×!<Ò!<Ø%×4Ñ4´e·k±kó"�Jô 1NØ%×4Ñ4ó1Ð-ð "ŸX™X×.Ñ.ØŸ	™	Ø4Ð5°Q°s°e¸4Ð@Ø!Ÿi™iØ+ð	 /ð �Fó $à!Ÿ™×,Ñ,×3Ñ3°D¸)ÀTÇYÁYÕOÙ",Ø)Ñ/Ø"ó#˜÷ Pð
 —I‘I×'Ñ'¨4¨&°°G°9Ð(=Ö>ñ- +S÷' ,ôV œS ×!>Ñ!>Ó?Ö@�Ø—	‘	×#Ñ#Ø<¸S¸EÐAtÐuxÐtyð  zDð  Eöó Að
 ×"×"Ð"¤s¨:£¸Ô':ä(¨×4‘��tØ—Y‘Y×%Ñ%¨UÐ%Ò3Ø!Ÿ[™[�FØ37×3M×3M¡ÐSV�Jà(,×(B×(B™È6È(ÐRWÐHXð ô
 —}‘}×(×(Ô-?Ó-AÀFÓ-JØ%7™
à%'˜
Ø—I‘I×'Ñ'Ø˜v˜hÐ&9¸*¸ÀRÈÀzÐQSÐTZ×T`ÑT`ÓTbÐScÐchÐisÐhtÐtvÐwô÷ 4ð —Y‘Y×%Ñ%¨U°Q©YÐ%Ò7Ø×8Ñ8¸¿y¹yÔI÷ 8Ò7ñ!  5ð( —‘×!Ñ!¬¨Z«Ð!Ò9Ø×.Ñ.¨t¯y©yÔ9Ø—	‘	× Ñ  ×!3Ñ!3Ô4Ø—	‘	× Ñ  §¡Ô,Ø—	‘	× Ñ  §¡Ô.Ø—	‘	× Ñ  §¡Ô-÷ :ô  (Ð(@¬)°JÓ*?Ð(@×A‘��tØ—Y‘Y×%Ñ%¨U°Q©YÐ%Ò7à26×2KÑ2KØŸ	™	ñ3ç‘e“gò3Ñ.˜	 ;ð !¤3 z£?°QÑ#6Ó6Ø(2°5¸1±9Ñ(=˜IØ/3×/HÑ/HØ )§¡ñ0à'ñ0)Ð,ô *7×)EÑ)EÀiÓ)P˜JÜ,3°I·O±OÀZÓ,P˜Mô 25°[ÐBRÔ1Sô+â1S¡I C¨ð !$ d¨]Ñ&:Ô :Ù1Sð (ñ +ð
 Ÿ	™	×+Ñ+Ü(Ø $× 8Ñ 8¸Ñ CØ#, +¨^¸I¸;ÀbÌÏÉ×I^ÑI^Ð_jÓIkÐHlÐlmÐ nóöñ!3÷ 8ð4 —‘×#Ñ# D×$:Ñ$:Ô;Ø× Ñ ×"ò9  Bð< �I‰I×Ñ˜T×/Ñ/Ô0Ø�I‰I×Ñ˜TŸZ™ZÔ(Ø�I‰I×Ñ˜TŸ\™\Ô*Ø�I‰I×Ñ˜TŸ[™[Ô)Ø�	‰	×Ñ˜×/Ñ/Ô0Ø×%×%Ø×"×" d×&:×&:à×-Ñ-Ð.Ð.AÐBˆGØ�I‰I×Ñð3à3:°)ð <ðð ð ñ ð ×6Ñ6×BÑBÔDØ×'×'Ø�I‰I×Ñ˜T×1Ñ1Ô2Ø×Ñ× Ñ Ô"Ø�
‰
×ÑÔØ�‰×ÑÔØ�‰×ÑÔØ×Ñ×$Ñ$Ô&Ø×Ñ×"Ñ"Õ$ùòs I÷r PÖOú÷I ,Ö+ú÷f 4Ö3ú÷ 8×7ú÷ :Ö9üó,+÷ 8Ö7úsw   Á6h&Âh&ÇH h=Ï0h+Ï<,h=ÓB%iÖi!×B0i4ÛBjÝ)jÝ>Ajè+
h:è5h=è=
ié
i	é!
i1	é4
jêjê
j	c                ó¶  • / nU R                  5       (       Ga@  / nU R                  SU/ 5        U GH$  n[        U[        5      (       a  UR	                  [        U5      5        M5  [        U[        5      (       aX  [        R                  R                  R                  UR                  U R                  S9nUR	                  [        U5      5        M¢  [        U[        R                  5      (       aO  [        R                  R                  R                  UU R                  S9nUR	                  [        U5      5        GM  [        S[!        U5       35      e   U$ )Nr°   ©rI  z!Unsupported numel argument type: )r�  Úadd_numel_to_call_argsrŠ   rî  rØ  rÔ   rl   rL   rP  rH  Úoptimization_hint_with_overrideÚ
inner_exprrI  r‹   rŒ   r{  rA  )r  rÂ  Ú
numel_argsr¡  Úhints        r†   Úkernel_benchmark_extra_argsÚ(TritonKernel.kernel_benchmark_extra_argsÌ  s  € ØˆØ×Ñ×!Ò!Ø+-ˆJØ×'Ñ'¨¨J¸Ô;Ü!�Ü˜c¤3×'Ñ'Ø—K‘K¤ C£Ö)Ü ¤_×5Ñ5ÜŸ7™7×+Ñ+×KÑKØŸ™Ø&*×&8Ñ&8ð Lð �Dð —K‘K¤ D£	Ö*Ü ¤U§Z¡Z×0Ñ0ÜŸ7™7×+Ñ+×KÑKØØ&*×&8Ñ&8ð Lð �Dð —K‘K¤ D£	×*ä$Ð'HÌÈcËÈÐ%TÓUÐUñ! "ð" ˆrˆ   c                óN  • [        5       nU R                  R                  5       u  p4pVUR                  / SQ5        UR	                  5          [
        R                  " 5       n/ n[        XE5       GH  u  pšS[        U5       3n[        R                  R                  U	5      nU(       a¾  [        R                  R                  R                  UR                  5       U R                  S9n[        R                  R                  R                  UR!                  5       U R                  S9nUR#                  U SU SU SUR%                  5        SUR'                  5        S3
5        GOU	[        R                  R(                  ;   aÓ  [        R                  R(                  U	   n[        R                  R                  R                  UR+                  5       U R                  S9n[        R                  R                  R                  UR-                  5       U R                  S9nUR#                  U SU SU SUR.                   SUR0                   S3
5        GO[3        U
[4        5      (       ae  [        R                  R                  R7                  U
R8                  U R                  S9nS	U
R:                  ;   a  S
nUR#                  U SU 35        O¥[3        U
[<        5      (       a‚  [        R                  R?                  5       n[        R                  R                  R7                  U
R                  U R                  S9nUR#                  U SU SU SU
R0                   S35        O[A        SU	 35      eURC                  U5        GM     URE                  U RG                  5       5        UR#                  SSRI                  U5       S35        SSS5        UR                  / SQ5        [        R                  R?                  5       nURJ                  nUR	                  5          UR#                  S[        R                  RL                  RO                  U5       S35        UR	                  5          UR#                  [        R                  RL                  RQ                  U5      5        SU 3nUR#                  U SU S35        UR#                  [S        [T        RV                  5       SU S35        SSS5        SSS5        UR                  / SQ5        UR	                  5          UR#                  S[        R                  RL                  RO                  U5       S35        UR	                  5          UR#                  [        R                  RL                  RQ                  U5      5        UR#                  S[S        [T        RV                  5       S35        SSS5        SSS5        UR                  / SQ5        UR	                  5          UR#                  S5        UR#                  S5        UR#                  S5        UR#                  S[        R                  R?                  5       RX                   S35        UR#                  SU 35        UR#                  S 5        UR#                  S!5        SSS5        U$ ! , (       d  f       GNó= f! , (       d  f       GNà= f! , (       d  f       GNê= f! , (       d  f       GN= f! , (       d  f       GN)= f! , (       d  f       U$ = f)"a�  
Generates Python code for benchmarking this Triton kernel.
- Creates example inputs (random tensors, constants, sizes).
- Runs the kernel on the current GPU/stream.
- Prints runtime (ms) and throughput (GB/s) using `num_gb`.
Args:
    num_gb (float): The number of gigabytes to use for throughput calculation.
Returns:
    IndentedBuffer: A buffer containing the generated Python benchmark code.
)r°   r°   zdef get_args():Úarg_r÷  z = rand_strided(r   z
, device='z	', dtype=r¡  rc  r   r  z = torch.zeros(z*Don't find the buffer or const tensor for r~  rè  N)Ú
r  zdef call(args):zwith rÕ  Ústreamz = get_raw_stream(z.run(*args, stream=)r  r  z def benchmark_all_configs(args):z.benchmark_all_configs(*args))r  r  zif __name__ == '__main__':z<from torch._inductor.runtime.benchmarking import benchmarkerr°   zargs = get_args()z7ms = benchmarker.benchmark(lambda: call(args), device='z
', rep=40)z	num_gb = zgb_per_s = num_gb / (ms / 1e3)z<print(f"{ms:.3f}ms    {num_gb:.3f}GB    {gb_per_s:.2f}GB/s"))-rV   rÂ  Úpython_argdefsrÛ  rž  r½  rD  r¤  r  rL   rP  Útry_get_bufferrH  Ú optimization_hints_with_overrideÚget_sizerI  Ú
get_striderh  Ú
get_devicer  Ú	constantsrÞ  ri  Údevicerƒ  rŠ   r\   rù  r’   rÈ   r^   r
  ÚKeyErrorrØ  Úextendrý  r   rþ   Ú
device_opsÚdevice_guardÚ
set_devicerÔ   rA   ÚKERNEL_NAMErA  )r  Únum_gbrk  Ú_argdefsÚ	call_argsr„  rW  Úname_cntÚ	var_namesÚarg_nameÚarg_sigÚvar_namer*  rÞ  ri  Úconst_tensorÚsymval_hintr
  rD  Úcurrent_devicerþ   Ústream_names                         r†   Úcodegen_kernel_benchmarkÚ%TritonKernel.codegen_kernel_benchmarkä  s  € ô  Ó!ˆØ,0¯I©I×,DÑ,DÓ,FÑ)ˆ˜Yà×ÑÒ5Ô6Ø�]‰]�_Ü —’Ó(ˆHØˆIÜ%(¨×%>Ñ!�Ø!¤$ x£.Ð!1Ð2�Ü—g‘g×,Ñ,¨XÓ6�ÞÜŸ7™7×+Ñ+×LÑLØŸ™›Ø&*×&8Ñ&8ð Mð �Dô ŸW™W×-Ñ-×NÑNØŸ™Ó(Ø&*×&8Ñ&8ð Oð �Fð ×$Ñ$Ø#˜*Ð$4°T°F¸"¸V¸HÀJÈsÏ~É~ÓO_ÐN`Ð`iÐjm×jwÑjwÓjyÐizÐz{Ð|öð ¤§¡×!2Ñ!2Ó2ä#$§7¡7×#4Ñ#4°XÑ#>�LÜŸ7™7×+Ñ+×LÑLØ$×)Ñ)Ó+Ø&*×&8Ñ&8ð Mð �Dô ŸW™W×-Ñ-×NÑNØ$×+Ñ+Ó-Ø&*×&8Ñ&8ð Oð �Fð ×$Ñ$Ø#˜*Ð$4°T°F¸"¸V¸HÀJÈ|×ObÑObÐNcÐclÐmy×mÑmð  mAð  ABð  Cöô   ¬×1Ñ1Ü"#§'¡'×"2Ñ"2×"RÑ"RØŸ™Ø&*×&8Ñ&8ð #Sð #�Kð %¨¯©Ó4Ø&'˜Ø×$Ñ$¨ z°°[°MÐ%BÕCÜ ¬×6Ñ6ÜŸW™W×@Ñ@ÓB�FÜŸG™G×,Ñ,×LÑLØŸ™°T×5GÑ5Gð Mð �Eð ×$Ñ$Ø#˜* O°E°7¸*ÀVÀHÈIÐV]×VcÑVcÐUdÐdeÐfõô #ØDÀXÀJÐOóð ð × Ñ  ×*ñk &?ðl ×Ñ˜T×=Ñ=Ó?Ô@Ø×Ñ˜w t§y¡y°Ó';Ð&<¸AÐ>Ô?÷u ðx 	×ÑÒ9Ô:ÜŸ™×<Ñ<Ó>ˆØ×$Ñ$ˆØ�]‰]�_Ø×Ñ˜u¤Q§W¡W×%7Ñ%7×%DÑ%DÀUÓ%KÐ$LÈAÐNÔOØ—‘•Ø× Ñ Ü—G‘G×&Ñ&×1Ñ1°%Ó8ôð !' u gÐ.�Ø× Ñ  K =Ð0BÀ5À'ÈÐ!KÔLØ× Ñ Üœ;×2Ñ2Ó3Ð4Ð4GÈÀ}ÐTUÐVô÷ !÷ ð 	×ÑÒJÔKØ�]‰]�_Ø×Ñ˜u¤Q§W¡W×%7Ñ%7×%DÑ%DÀUÓ%KÐ$LÈAÐNÔOØ—‘•Ø× Ñ Ü—G‘G×&Ñ&×1Ñ1°%Ó8ôð × Ñ Øœc¤+×"9Ñ"9Ó:Ð;Ð;XÐYô÷	 !÷ ð 	×ÑÒDÔEØ�]‰]�_Ø×ÑØNôð ×Ñ˜RÔ à×ÑÐ0Ô1Ø×ÑØIÌ!Ï'É'×JmÑJmÓJo×JtÑJtÐIuÐuð  Aôð ×Ñ˜y¨¨Ð1Ô2Ø×ÑÐ=Ô>Ø×ÑØNô÷ ð  ˆ÷Q Ž_ú÷B !–ú÷ Ž_ú÷  !–ú÷ Ž_ú÷ Œ_ð  ˆúsf   ÁNZ;Ð$A[Ñ1B[Ó4[Ô'A\Õ4A%[1×\ØB%\Ú;
[
Û
[	Û[Û
[.Û1
\ 	Û;\Ü
\Ü
\$c                óš   • [         R                  " SR                  [        R                  R
                  R                  S5      5      5      $ )Nzl
            from torch._dynamo.testing import rand_strided
            {}
            import torch
        Úget_raw_stream)ÚtextwrapÚdedentrÃ  rL   rP  r  Úimport_get_raw_stream_asr  s    r†   Úimports_for_benchmark_kernelÚ)TritonKernel.imports_for_benchmark_kernel]  s:   € Ü�Šð÷ ‰F”1—7‘7×%Ñ%×>Ñ>Ð?OÓPÓQó
ð 	
rˆ   c                óº   • U R                   (       a  gU R                  (       a  gU R                  (       a  U R                  (       d   egU R                  (       a  gg)Nr  rQ  r  r7  Ú	pointwise)r  rQ  r  rY  r  s    r†   Ú_get_heuristicÚTritonKernel._get_heuristicf  sD   € Ø××Ø!Ø×'×'Ø*Ø×&×&Ø×(×(Ð(Ð(Ø)Ø×"×"ØØrˆ   c                ój  • [        5       nUR                  S5         SS KnUR                  S5        [	        5       =n(       a  UR                  U5        UR                  S5        [        R                  R                  (       a  UR                  S5        UR                  5       $ ! [         a     Nf = f)NzP
            import triton
            import triton.language as tl
            r   zU
               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')
                )
rV   rz  Útriton.language.extra.tlxÚImportErrorr¶   rh  r"   rµ   Úproton_profilingr‚  )rŸ   Úimportsrµ   Ú	attr_descs       r†   Úgen_common_triton_importsÚ&TritonKernel.gen_common_triton_importsr  sª   € ô !Ó"ˆØ�‰ðô	
ð		Û,à�N‰Nðôô 3Ó4Ð4ˆ9Õ4Ø×Ñ˜iÔ(à�‰ðô	
ô �=‰=×)×)Ø�N‰Nðôð ×ÑÓ!Ð!øô+ ó 	Ùð	ús   �B% Â%
B2Â1B2c                ó€   • [         R                  (       + U R                  5       [         R                  R                  S.$ )N)Úenable_fp_fusionÚ
launch_pdlÚdisable_ftz)r"   Úemulate_precision_castsr0  rö  r5  r©  s    r†   Útriton_meta_commonÚTritonKernel.triton_meta_common›  s5   € ô %+×$BÑ$BÔ BØ×1Ñ1Ó3Ü!×0Ñ0×<Ñ<ñ
ð 	
rˆ   c                ó¬  • [         R                  R                  R                  5       [        R
                  [        R                  [        R                  R                  [        R                  [        R                  [        R                  [        R                  [        R                  [        R                  R                  [        R                  R                  [        R                  R                   [        R"                  [        R$                  R&                  [        R                  R(                  S.n[        R*                  (       a  [         R,                  " 5       US'   [         R.                  R0                  b  SUS'   [        R2                  " 5       (       a  SUS'   [        R4                  (       aL  [        R4                  US'   [        R6                  US'   [        R8                  US'   [        R:                  US	'   [        R<                  (       a9  [        R<                  US
'   [        R>                  US'   [        R@                  US'   U$ )N)Úbackend_hashÚassert_indirect_indexingÚautotune_local_cacheÚautotune_pointwiseÚautotune_remote_cacheÚforce_disable_cachesÚdynamic_scale_rblockÚmax_autotuneÚmax_autotune_pointwiseÚmin_split_scan_rblockÚspill_thresholdÚstore_cubinr]  r^  Ú&mix_order_reduction_allow_multi_stagesÚ$are_deterministic_algorithms_enabledTÚis_hiprï  Úprofile_bandwidthÚprofile_bandwidth_regexÚprofile_bandwidth_outputÚ/profile_bandwidth_with_do_bench_using_profilingÚcoordinate_descent_tuningÚ coordinate_descent_search_radiusÚ'coordinate_descent_check_all_directions)!rœ   r%   Ú_tritonÚtriton_hash_with_backendr"   r;  r<  rµ   r=  r>  r?  r@  rA  rB  rC  rD  rE  r]  rG  r^  rF  Ú*write_are_deterministic_algorithms_enabledrG  rð  rñ  rï  rI  rJ  rK  rL  rM  rN  rO  )rŸ   Úinductor_metas     r†   Úinductor_meta_commonÚ!TritonKernel.inductor_meta_common£  s³  € ô "ŸK™K×/Ñ/×HÑHÓJÜ(.×(GÑ(GÜ$*×$?Ñ$?Ü"(§-¡-×"BÑ"BÜ%+×%AÑ%AÜ$*×$?Ñ$?Ü$*×$?Ñ$?Ü"×/Ñ/Ü&,×&CÑ&CÜ%+§]¡]×%HÑ%HÜ%Ÿ}™}×<Ñ<Ü!Ÿ=™=×4Ñ4Ü#×1Ñ1Ü.4×.AÑ.A×.`Ñ.`Ü6<·m±m×6jÑ6jñ
ˆô$ ×<×<ä×:Ò:Ó<ð Ð@ÑAô �=‰=×ÑÑ(Ø&*ˆM˜(Ñ#Ü×Ò×ÑØ)-ˆM˜+Ñ&Ü×#×#Ü17×1IÑ1IˆMÐ-Ñ.Ü7=×7UÑ7UˆMÐ3Ñ4Ü8>×8WÑ8WˆMÐ4Ñ5ä×FÑFð ÐKÑLô ×+×+ä×0Ñ0ð Ð5Ñ6ô ×7Ñ7ð Ð<Ñ=ô ×>Ñ>ð ÐCÑDð Ðrˆ   c                óx  ^-^.• [        5       n0 nU R                  R                  5        H‘  u  pE[        U5      (       a  U R                  (       d  M(  [
        R                  R                  R                  U5      n[        U[        [        R                  45      (       d  SnO[        [        U5      5      nXsU'   M“     Uc¤  UR                  U R                  5       5        [
        R                  R!                  5       R"                  nUS:X  a  UR                  S5        OUR                  S5        [$        R&                  (       a  UR                  U R)                  5       5        U R*                  R-                  5       u  m-n	m.n	[/        T.5       H¨  u  p«[        U[0        5      (       d  M  [3        [        R4                  UR6                  5      nU[
        R                  R                  R8                  ;   d  Mj  [1        UR:                  [
        R                  R                  R8                  U   5      T.U
'   Mª     [=        5       nU R>                   GH'  nXàR*                  R@                  ;   a(  URC                  U R*                  R@                  U   5        XàR*                  RD                  ;   am  U[
        R                  RF                  ;  aO  XàRF                  ;  a@  URC                  [3        [H        U R*                  RD                  U   5      RJ                  5        XàR*                  RL                  ;   d  Mæ  U R*                  RL                  U   n[        U[N        5      (       a   eURC                  U5        GM*     [Q        T-T.5       HX  u  nn[        U[R        5      (       d  M  URT                  [V        RX                  :X  d  M=  URC                  UR:                  5        MZ     [[        U5      nU R]                  5        H[  n[1        UR^                   S3UR`                  5      nT.Rc                  U5        T-Rc                  [e        UR:                  5      5        M]     U-U.4S jnU Rf                   HY  nURh                  (       a  U Rj                  (       a  M'  URl                  c  M6  U" UR^                  Ro                  5        S35        M[     U Rp                  (       a  U" S	5        U Rr                  (       a  U" S
5        U" S5        [u        T.U Rv                  T-S9nU[x        Rz                  " [
        R                  R!                  5       5      0 [|        R~                  R$                  R€                  R‚                  =(       a7    S[…        U R†                  5      ;   =(       d    S[…        U Rˆ                  5      ;   S.U R‹                  5       EnU Rp                  (       a  SUS'   [
        R                  RŒ                  =(       d    [
        R                  RŽ                  nU R‘                  5       R’                  [•        U R–                  5      […        [˜        Rš                  5      UUU Rœ                  U Rž                  U R                   U R¢                  U R¤                  S.
U R§                  5       EnU Rr                  (       a  U R¨                  US
'   [$        Rª                  (       d  [$        R¬                  R®                  (       a"  U R°                  =(       d    U R²                  US'   [
        R´                  R¶                  Ri                  5       =(       a    U Rj                  (       + nU R¸                  n[»        U R¼                  5      S:H  nU(       Gañ  U(       Gaé  U R¶                  R¿                  U R¼                  5      nURÀ                  RÂ                  RÄ                  S   nURÆ                  nUb+  SU;   a%  SU;   a  US   [É        US   S5      -  nU[Ê        :¬  nO,U R¶                  RÍ                  U5      [Î        RÐ                  :H  nURÒ                  RÂ                  RÔ                  n URÀ                  RÂ                  RÔ                  n![
        R                  R                  R×                  U 5      [É        [
        R                  R                  R×                  U!5      S5      -  n"U"S:¼  a˜  U(       a‘  [
        R                  R                  RÙ                  U R¶                  RÚ                  S5      (       aN  [
        R                  R                  RÝ                  U R¶                  RÚ                  S5      (       a  US::  a  SUS'   U R¸                  (       a  U R¸                  US'   U RÞ                  (       a  U RÞ                  US'   U Rp                  (       a  U Rj                  US'   Sn#[$        R&                  (       d  [$        Rà                  (       a  U Rã                  5       S -  n#U#b  U#US!'   [$        R&                  (       a  U Rå                  5       n$U$b  U$US"'   [ç        T.5      /US#'   [é        T.5       H  n%SUS$   T.U%   R:                  '   M     UU lu        UU lv        U Rï                  U R†                  5        U Rñ                  5         U Ró                  U R†                  5        U Rô                   H%  n&UR÷                  S%5        UR                  U&5        M'     U Rø                  (       a5  S&U Rû                  5        S'U Rø                  R$                  < S(U< S)U< S*3	n'O¸U R                  (       aI  U R¶                  RÍ                  U R¸                  5      n(S&U Rû                  5        S+U< S,U( S(U< S)U< S*3n'O^S%n)[»        U5      S:X  a  [»        [ý        T.5      5      S-:X  a  S.n)OS/n)S&U Rû                  5        S+U< S0U) S1U< S)U< S2U Rþ                   S*3n'UR                  U'5        U=(       d    […        [˜        GR                   5      n*UR÷                  S3U* S4S0GR                  S5 T- 5       5       S635        UGR                  5          [$        R€                  GR                  (       a  UR÷                  S7U* S835        U GR	                  U5        U R*                  GR                  5        H  u  n+n,UR÷                  U+ S9U, 35        M     UR                  U R†                  5        [$        R€                  GR                  (       a  UR÷                  S:U* S835        SSS5        [$        R&                  (       a!  UR                  U GR                  U#5      5        UGR                  5       $ ! , (       d  f       NU= f);z‘
Convert the TritonKernel from Inductor SIMD IR to triton code, including inductor triton heuristics, imports,
metadata, and benchmarking infra.
i    NÚcpuz"triton_helpers.set_driver_to_cpu()z"triton_helpers.set_driver_to_gpu()r  c                óŠ   >• [        5       (       a  TR                  [        U 5      5        TR                  [        U SS95        g )NT)Úis_constexpr)rG   rØ  rR   rP   )r  Úargdefsr„  s    €€r†   Úadd_constexpr_argÚ6TritonKernel.codegen_kernel.<locals>.add_constexpr_arg:  s2   ø€ ä-×/Ñ/Ø× Ñ ¤¨hÓ!7Ô8Ø�N‰Nœ7 8¸$Ñ?Õ@rˆ   r¾   rZ  ÚRSPLIT_SIZEÚ
NUM_STAGES)Ú
size_dtyperZ  ztl.dot)r„  r
  r	  Únative_matmulTÚlaunch_cooperative_grid)
Ú	grid_typerN  Úkernel_nameÚmutated_arg_namesÚoptimize_memrV  rŠ  Únum_loadÚ	num_storeÚnum_reductionÚ"has_loadstore_with_contiguous_rdimr!   r   rN  Úr0_rN   gÍÌÌÌÌÌô?i €  r‘  é
   Úadd_persistent_rblockÚtiling_scoresr   r  g    eÍÍAÚkernel_num_gbÚkernel_flopÚconfigsr	  r°   z$
                @triton_heuristics.z(
                    config=zI,
                    filename=__file__,
                    triton_meta=z$,
                    inductor_meta=z;
                )
                @triton.jit
            z!(
                    size_hints=z%,
                    reduction_hint=r‰  ztile_hint=TileHint.SQUARE,ztile_hint=TileHint.DEFAULT,r   zH
                    filename=__file__,
                    triton_meta=z*,
                    min_elem_per_thread=zdef r)  c              3  ó@   #   • U  H  oR                  5       v •  M     g 7frš   )Ú	full_namer:  s     r†   r  Ú.TritonKernel.codegen_kernel.<locals>.<genexpr>  s   é € Ð*JÂ'¸Q¯;©;¯=¨=Â'ùrÈ  rn  zpl.enter_scope("rD  r  zpl.exit_scope(")ˆrV   rZ  r¼  rB   rY  rL   rP  rH  r  rŠ   rî  r‹   r�   r5   rz  r0  r
  rA  r"   Úbenchmark_kernelr$  rÂ  r  r–  r\   r
   rô   r’   Úinv_precomputed_replacementsrÈ   r   Ú	mutationsÚinput_buffersr„  rˆ  Úremoved_buffersrW   Ú
inner_nameÚoutput_buffersr[   r¤  r^   Ú	zero_moder_   ÚZERO_ON_CALLr!  rÒ   rÓ   r  rØ  rP   re  r\  r  rÖ   rõ   rQ  rÚ  rk   ræ  r0   rl  rœ   r4  rµ   r`  rÔ   re  r5  r7  Úis_inferenceÚis_backwardÚ_get_grid_typer›   ÚsetrN  rA   ÚDESCRIPTIVE_NAMErV  rŠ  rf  rg  rh  rT  Úrsplit_sizer]  rG  r^  rU  ri  rÅ   r[  rm  r�   rW  Úmemory_statsÚ
persistentÚmemoryrg  Úcount_per_threadr{   ÚINNER_REDUCTION_RATIO_THRESHOLDÚget_reduction_hintr1   ÚINNERÚloopedÚbytesÚoptimization_hintÚstatically_known_leqÚreduction_numelÚstatically_known_gtr   rI  Úestimate_kernel_num_bytesÚestimate_flopsrg   rh   rO  rS  rÐ  rô  r[  r5  rh  r  r(  ri   rÚ  r  r   rž  r-  Úcodegen_static_numelsÚaliasesr  r‚  )/r  rÈ   rV  Ú
size_hintsrÓ   r  Ú
numel_hintÚ	size_hintÚdevice_typerW  rU  r¡  r£  Úmutated_argsÚmutationÚmutation_argÚargnamerá   Úsizeargr[  Útriton_meta_signaturerO  re  rS  Ú
looped_redrm  Ú	two_d_redrƒ  Ú	dim_statsÚmem_ops_per_threadÚr_coalesce_ratioÚcontiguous_redÚ
looped_memÚpersistent_memÚsaved_bytes_ratior  ÚflopsÚarg_numrv  Úheuristics_lineÚreduction_hintÚ	tile_hintrc  ÚoldÚnewrZ  r„  s/                                                @@r†   Úcodegen_kernelÚTritonKernel.codegen_kernelÓ  sõ
  ù€ ô Óˆàˆ
Ø!Ÿ[™[×.Ñ.Ö0‰MˆFÜ" 6×*Ñ*°4×3H×3HÙäŸ™×)Ñ)×7Ñ7¸Ó>ˆJÜ˜j¬3´·±Ð*>×?Ñ?ð !‘	ä+¬C°
«OÓ<�	Ø!*�vÓñ) 1ð, ‰<Ø�K‰K˜×6Ñ6Ó8Ô9ÜŸ'™'×=Ñ=Ó?×DÑDˆKØ˜eÓ#Ø—‘Ð@ÕAà—‘Ð@ÔAä×&×&Ø—‘˜D×=Ñ=Ó?Ô@à#'§9¡9×#;Ñ#;Ó#=Ñ ˆ��I˜qä 	Ö*‰FˆAÜ˜#œw×'Ó'ô œeŸl™l¨C¯H©HÓ5�ØœQŸW™W×-Ñ-×JÑJÕJÜ#*ØŸ™¤!§'¡'×"2Ñ"2×"OÑ"OÐPVÑ"Wó$�I˜a“Lñ +ô )3«ˆØŸ�ˆHØŸ9™9×2Ñ2Ó2Ø× Ñ  §¡×!8Ñ!8¸Ñ!BÔCàŸI™I×5Ñ5Ó5Ø¤A§G¡G×$;Ñ$;Ó;Ø×$8Ñ$8Ó8à× Ñ Üœ¨¯©×)BÑ)BÀ8Ñ)LÓM×XÑXôð Ÿ9™9×3Ñ3Õ3Ø#Ÿy™y×7Ñ7¸ÑA�Ü% l´J×?Ñ?Ð?Ð?Ø× Ñ  ×.ñ 'ô6   ¨Ö3‰LˆG�Sä˜3¤×-Ó-Ø—M‘MÔ%6×%CÑ%CÕCà× Ñ  §¡Ö.ñ 4ô ˜lÓ+ˆà×+Ñ+Ö-ˆDÜ §¡ ¨UÐ3°T·Z±ZÓ@ˆGØ×Ñ˜WÔ%Ø�N‰Nœ7 7§<¡<Ó0Ö1ñ .ö	Að ×$Ô$ˆDØ× ×  T×%>×%>áØ�‰Ñ&Ùá §¡×!2Ñ!2Ó!4Ð 5°UÐ;Ö<ñ %ð ×%×%Ù˜hÔ'à×#×#Ù˜mÔ,Ù˜lÔ+ä 1Ø $×"2Ñ"2¸Gñ!
Ðð /Ü&×-Ò-¬a¯g©g×.QÑ.QÓ.SÓTØä—‘×&Ñ&×-Ñ-×;Ñ;÷ RØ¤ T§Y¡Y£Ñ/×P°8¼sÀ4Ç<Á<Ó?PÑ3Pñ	'
ð ×%Ñ%Ó'ð	'
ˆð ×%×%ð 6:ˆKÐ1Ñ2ô
 —w‘w×+Ñ+×B¬q¯w©w×/BÑ/Bˆð ×,Ñ,Ó.×7Ñ7ä! $×"5Ñ"5Ó6Üœ{×;Ñ;Ó<Ø!-Ø(ØŸ™Ø $× 5Ñ 5ØŸ™ØŸ™Ø!×/Ñ/ñ
ð ×'Ñ'Ó)ð
ˆð ×#×#Ø+/×+;Ñ+;ˆM˜-Ñ(ä××¤6×#6Ñ#6×#U×#Uà×2Ñ2÷ 7Ø×6Ñ6ð Ð>Ñ?ô —X‘X×&Ñ&×3Ñ3Ó5×W¸d×>WÑ>WÔ:Wˆ
Ø×*Ñ*ˆÜ˜Ÿ™Ó$¨Ñ)ˆ	ßŸ)ØŸ=™=×5Ñ5°d·k±kÓBˆLØ$×/Ñ/×6Ñ6×:Ñ:¸1Ñ=ˆIØ!*×!;Ñ!;Ðð Ñ)Ø˜=Ó(Ø˜]Ó*ð $1°Ñ#7¼#¸mÈCÑ>PÐRSÓ:TÑ#TÐ Ø!1Ô5TÑ!T‘ð —M‘M×4Ñ4°]ÓCÜ$×*Ñ*ñ+ð ð
 &×,Ñ,×3Ñ3×9Ñ9ˆJØ)×4Ñ4×;Ñ;×AÑAˆNä !§¡× 0Ñ 0× BÑ BÀ:Ó NÔQTÜ—‘× Ñ ×2Ñ2°>ÓBØóRñ !Ðð " SÓ(Þ"ä—G‘G×$Ñ$×9Ñ9Ø—M‘M×1Ñ1°5÷ñ ô —G‘G×$Ñ$×8Ñ8Ø—M‘M×1Ñ1°4÷ñ ð '¨"Ó,à9=�Ð5Ñ6à××Ø-1×-?Ñ-?ˆM˜/Ñ*à×#×#Ø37×3KÑ3KˆMÐ/Ñ0à×%×%Ø48×4MÑ4MˆMÐ0Ñ1àˆÜ×"×"¤f×&>×&>Ø×3Ñ3Ó5¸Ñ;ˆFØÑ!Ø17�˜oÑ.Ü×"×"Ø×'Ñ'Ó)ˆEØÑ Ø/4�˜mÑ,ä"+¨IÓ"6Ð!7ˆ�IÑô +¨9Ö5ˆGØ@AˆK˜Ñ$ Y¨wÑ%7×%<Ñ%<Ó=ñ 6ð 'ˆÔØ*ˆÔà×Ñ˜dŸi™iÔ(Ø×ÑÔØ×Ñ˜Ÿ™Ô#à×+Ô+ˆFØ�N‰N˜2ÔØ�K‰K˜Öñ ,ð ××ð#$Ø$(×$7Ñ$7Ó$9Ð#:ð ;Ø ×-Ñ-×4Ñ4Ñ7ð 8!à!,¡ð 0#Ø#0Ñ"3ð 4ð‰Oð ×"×"Ø!Ÿ]™]×=Ñ=¸d×>PÑ>PÓQˆNð#$Ø$(×$7Ñ$7Ó$9Ð#:ð ; Ø *™~ð .$Ø$2Ð#3ð 4!à!,¡ð 0#Ø#0Ñ"3ð 4ð	‰Oð ˆIÜ�:‹ !Ó#äÔ/°	Ó:Ó;¸qÓ@à <‘Ià =�Ið#$Ø$(×$7Ñ$7Ó$9Ð#:ð ; Ø *™~¨R°	¨{ð ;!à!,¡ð 0#Ø#0Ñ"3ð 4)Ø)-×)AÑ)AÐ(Bð Cð	ˆOð 	�‰�OÔ$Ø×:œc¤+×"9Ò"9Ó:ˆØ�‰Ø�;�-˜q §¢Ñ*JÁ'Ó*JÓ!JÐ KÈ2ÐNô	
ð �[Š[�]Ü�}‰}×-×-Ð-Ø—‘Ð!1°+°¸bÐAÔBØ×&Ò& tÔ,Ø ŸI™I×-Ò-Ö/‘��SØ—‘ #  c¨#¨Ð/Ö0ñ 0à�K‰K˜Ÿ	™	Ô"Ü�}‰}×-×-Ð-Ø—‘ °°¸RÐ@ÔA÷ ô ×"×"Ø�K‰K˜×5Ò5°fÓ=Ô>à�}Š}‹Ð÷ �]ús   ò	Cv+ö+
v9c                óÄ  • [         R                  R                  R                  U 5      n [	        U [
        R                  [        45      (       a  [        U 5      n[        U5      nU$ Sn[         R                  R                  R                  X5      (       dI  US:”  a  [        SU  35      eUS-  n[         R                  R                  R                  X5      (       d  MI  U$ )NrN   i @  z!Failed to find static RBLOCK for r!   )rL   rP  rH  ÚsimplifyrŠ   r‹   r�   rî  r5   r�  r{  )rÉ  rÃ  s     r†   r  Ú#TritonKernel._get_persistent_RBLOCK  sµ   € ä—‘×!Ñ!×*Ñ*¨6Ó2ˆÜ�fœuŸ}™}¬cÐ2×3Ñ3Ü�f“+ˆCÜ! #Ó&ˆCð ˆ
ð ˆCÜ—g‘g×&Ñ&×;Ñ;¸F×HÑHØ˜“?Ü$Ð'HÈÈÐ%QÓRÐRØ�q‘�ô —g‘g×&Ñ&×;Ñ;¸F×HÓHð
 ˆJrˆ   c                óP   •  [         R                  U 5        g! [         a     gf = f)NTF)r  r  r{  )rÉ  s    r†   Úhas_persistent_RBLOCKÚ"TritonKernel.has_persistent_RBLOCK'  s*   € ð	Ü×/Ñ/°Ô7ØøÜó 	Ùð	ús   ‚ ˜
%¤%c                óR  • S
S jnU R                    GH‘  nUR                  (       a  U R                  (       ai  [        R                  R
                  R                  UR                  5      nU" U5      (       a)  UR                  UR                   S[        U5       35        UR                  (       a¹  U R                  (       a¨  U R                  (       a1  U R                  U R                  UR                  5      5      nSU S3nO8U R                  UR                  5      nU R                   (       a  [#        US5      nUR                  UR                  R%                  5        SU 35        UR                  S:X  d  GMl  U R&                  (       d  GM€  UR                  S5        GM”     g	)ay  
We get a small speedup from hard coding numels if they are static.

This code stomps on the passed-in values by writing an constant to the top of the kernel.

In a kernel like:
def KERNEL_NAME(in_ptr0, in_ptr1, out_ptr2, xnumel, r0_numel, XBLOCK : tl.constexpr, R0_BLOCK : tl.constexpr):

We would add
xnumel = 4096
r0_numel = 768

After the signature, before the kernel code, if we decided to make these static. As its hardcoded, it becomes
a better signal to triton on how to unroll and do some static indexing. So, it's not so much that downstream
knows that its a static numel, as that you just plop a constant into the kernel.
c                óB   • [        U [        R                  [        45      $ rš   )rŠ   r‹   r�   rî  r‘   s    r†   Úis_static_integerÚ=TritonKernel.codegen_static_numels.<locals>.is_static_integerA  s   € Ü˜d¤U§]¡]´CÐ$8Ó9Ð9rˆ   znumel = z*triton_helpers.constexpr_next_power_of_2((z + RSPLIT - 1) // RSPLIT)r  zBLOCK: tl.constexpr = rN  zXBLOCK: tl.constexpr = 1N)r’   rí   r¦   r¥   )re  r\  rY  rL   rP  rH  r±  r  rh  rÓ   rî  r  rQ  r6  r   r  rE  r{   rõ   rV  )r  rV  r¸  rá   Úsimplified_tree_numelr  rÃ  s          r†   r’  Ú"TritonKernel.codegen_static_numels/  s,  € ô$	:ð ×$Õ$ˆDØ×$×$¨×(=×(=Ü()¯©×(8Ñ(8×(AÑ(AÀ$Ç*Á*Ó(MÐ%Ù$Ð%:×;Ñ;Ø—N‘N d§k¡k ]°(¼3Ð?TÓ;UÐ:VÐ#WÔXà× ×  T×%>×%>Ø×-×-Ø ŸJ™J t×';Ñ';¸D¿J¹JÓ'GÓH�EØFÀuÀgÐMfÐg‘Cà×5Ñ5°d·j±jÓA�CØ×,×,ä! # r›l˜à—‘ $§+¡+×"3Ñ"3Ó"5Ð!6Ð6LÈSÈEÐRÔSà�{‰{˜cÖ! d§m§m¢mØ—‘Ð9×:ò' %rˆ   c                óŠ  • [        U R                   Vs/ s H  n[        UR                  (       + 5      PM     sn5      nU R                  (       a  US:X  d   e[
        R                  $ U R                  (       a  US:X  d   e[
        R                  $ US:X  a  [
        R                  $ US:X  aN  [        [        U R                  U R                  5      5      (       a  [
        R                  $ [
        R                  $ US:X  a1  U R                  (       a  [
        R                   $ [
        R"                  $ [%        SU 35      es  snf )NrN   r!   r   z"Unsupported number of dimensions: )rÍ  re  rî  r\  rÚ  r-   ÚMixOrderReductionGridrQ  ÚCooperativeReductionGridÚGrid1Drz   r"  r  ÚGrid2DWithYZOverflowÚGrid2DrE  ÚBatchMatmulGrid3DÚGrid3Dr{  )r  rá   r¡  s      r†   r  ÚTritonKernel._get_grid_typeY  s  € Ü¸×8HÒ8HÓIÒ8H°”˜×*Ñ*Ô*Ö+Ñ8HÑIÓJˆØ×#×#Ø˜“6ˆM�6Ü$×:Ñ:Ð:Ø×'×'Ø˜“6ˆM�6Ü$×=Ñ=Ð=Ø�!‹VÜ$×+Ñ+Ð+Ø�!‹VÜ”3�t×2Ñ2°D×4DÑ4DÓE×FÑFÜ(×=Ñ=Ð=Ü$×+Ñ+Ð+Ø�!‹VØ×$×$Ü(×:Ñ:Ð:Ü$×+Ñ+Ð+ÜÐ=¸a¸SÐAÓBÐBùò# Js   ”$E c                ó¦  • U R                    HÁ  n[        UR                  [        R                  [        R
                  45      (       a  UR                  nO)[        R                  R                  R                  X5      nUR                  (       a  U R                  (       d  M–  UR                  U5        UR                  [        U5      5        MÃ     g rš   )re  rŠ   r  r‹   r�   rô   rL   rP  Úwrapper_codeÚgenerate_numel_exprr\  rY  rØ  rA  )r  rÈ   r  Ú	arg_typesrá   r’   s         r†   rø  Ú#TritonKernel.add_numel_to_call_argsm  s‡   € à×$Ô$ˆDÜ˜$Ÿ*™*¤u§}¡}´e·l±lÐ&C×DÑDØ—z‘z‘ä—w‘w×+Ñ+×?Ñ?ÀÓK�à×$×$¨×(=×(=Ñ(=Ø× Ñ  Ô&Ø× Ñ ¤ d£Ö,ò %rˆ   c           	     ó’  • [         R                  R                  nUR                  5         U R                  R                  5       u  pVpWU R                  XU5        U R                  R                   H  nUR                  U5        M     UR                  UUSUU R                  U R                  S9  U(       a  U R                  5         g g )NT)rµ   rÈ  rO  rS  )rL   rP  rÆ  Úwrite_triton_header_oncerÂ  r  rø  Úworkspace_argsÚgenerate_workspace_allocationÚgenerate_kernel_callrO  rS  Údeallocate_workspaces)	r  rÈ   r[  Údeallocate_wsÚwrapperrW  r  rÈ  Úwss	            r†   Úcall_kernelÚTritonKernel.call_kernely  s±   € ô —'‘'×&Ñ&ˆØ×(Ñ(Ô*Ø%)§Y¡Y×%=Ñ%=Ó%?Ñ"ˆ�aØ×#Ñ# D°YÔ?à—)‘)×*Ô*ˆBØ×1Ñ1°"Ö5ñ +ð 	×$Ñ$ØØØØØ×(Ñ(Ø×,Ñ,ð 	%ñ 	
ö Ø×&Ñ&Õ(ð rˆ   c                ó–  • [         R                  R                  nU R                  R	                  5       u  p#pB[        X45       Hƒ  u  pV[        U[        5      (       d  M  [         R                  R                  (       a  UR                  SU SU S35        MU  SU S3nUR                  U5        SU S3nUR                  U5        M…     g )Nz:AOTI_TORCH_ERROR_CODE_CHECK(aoti_torch_check_inf_and_nan("z", z));zassert not z.isnan().any().item()z.isinf().any().item())
rL   rP  rÆ  rÂ  r  r¤  rŠ   r]   Úcpp_wrapperrh  )r  rÑ  rW  r  Úarg_signaturesr¡  Úarg_signaturer%  s           r†   Úcodegen_nan_checkÚTritonKernel.codegen_nan_check�  s·   € Ü—'‘'×&Ñ&ˆØ*.¯)©)×*BÑ*BÓ*DÑ'ˆ�nÜ"% iÖ"@ÑˆCÜ˜-¬×3Ó3Ü—7‘7×&×&Ø×%Ñ%ØTÐUXÐTYÐY\Ð]`Ð\aÐadÐeöð )¨¨Ð-BÐC�DØ×%Ñ% dÔ+Ø(¨¨Ð-BÐC�DØ×%Ñ% dÖ+ò #Arˆ   c                ó   • [        U0 UD6$ rš   )r–  )r  rÂ  r   s      r†   ra  ÚTritonKernel.create_cse_varŸ  s   € Ü  $Ð1¨&Ñ1Ð1rˆ   c                ó\  • UR                    SU R                  U R                  UR                  5      5       3nUR                  R
                  (       d!  U R                  (       a,  UR                  S:X  a  U R                  R                  U5        g U R                  R                  U5        g )Nr  rN  )rÈ   r6  r   r’   Úrootrƒ  rÚ  rÓ   rÙ  rh  re  )r  rå  r%  s      r†   Úcodegen_iteration_ranges_entryÚ+TritonKernel.codegen_iteration_ranges_entry¢  sy   € Ø—*‘*�˜S §¡¨D×,@Ñ,@ÀÇÁÓ,LÓ!MÐ NÐOˆð �:‰:×× $×":×":¸u¿|¹|ÈsÓ?RØ×Ñ×(Ñ(¨Õ.ð �I‰I×Ñ Õ%rˆ   c                ó¬  • UR                   c   eU R                  UR                   5      nU R                  (       a&  UR                   S:X  a  U R                  5       S:X  a  SnU R                  nUS:w  a  SU S3OSnU R
                  (       a'  U R                  (       a  UR                  (       a  U S3nSUR                  R                  5        S	U U 3$ )
Nr   r‰  r°   rŸ  r»  r¡  z + rsplit_startztl.arange(0, zBLOCK))
rÖ   Úindexing_size_strrE  rÑ   ræ  rQ  r  r\  rÓ   rõ   )r  rå  rÞ  ræ  r>  s        r†   r‡  Ú)TritonKernel.iteration_ranges_ranges_code­  sÇ   € Ø×ÑÑ+Ð+Ð+Ø×%Ñ% e×&6Ñ&6Ó7ˆð ×!×!Ø× Ñ  AÓ%Ø×'Ñ'Ó)¨QÓ.àˆDØ×&Ñ&ˆØ*5¸Ó*C�4˜�} AÑ&Èˆà×&×&Ø×)×)Ø×"×"à�x˜Ð/ˆFØ˜uŸ|™|×1Ñ1Ó3Ð4°F¸4¸&ÀÀÐIÐIrˆ   c                ó^   • U R                   nU R                  5       nS/U-  nSU SU SU S3$ )NrN   rå  r   r¡  )ræ  rÑ   )r  rå  rª  ræ  rà   rÞ  s         r†   Úiteration_ranges_scalar_codeÚ)TritonKernel.iteration_ranges_scalar_codeÂ  sC   € ð ×&Ñ&ˆØ×&Ñ&Ó(ˆØˆs�T‰zˆØ˜$˜˜r % ¨¨;¨-°qÐ9Ð9rˆ   c                óª  • UR                   c   eSUR                    S3nU R                  U5      (       a#  SU SUR                   S-    SUR                    S3nU R                  (       a,  U R                  5       S:X  a  S	SS
S.nSX1R                       S3nUR                  R                  X"5      nU R                  S:w  a  U SU R                   S3$ U$ )Nztl.program_id(r¡  r)  z + tl.program_id(rN   z) * tl.num_programs(r@  r‰  r!   r   )r   rN   r!   rŸ  r»  )rv  r  rE  rÑ   Ú	pid_cacherè  ræ  )r  rå  rÂ  Úreversed_pid_mapÚpids        r†   Úiteration_ranges_get_pidÚ%TritonKernel.iteration_ranges_get_pidÊ  sé   € Ø�~‰~Ñ)Ð)Ð)Ø˜uŸ~™~Ð.¨aÐ0ˆð ×&Ñ& u×-Ñ-ð �c�UÐ+¨E¯N©N¸QÑ,>Ð+?Ð?SÐTY×TbÑTbÐScÐceÐfˆCð
 × ×  T×%<Ñ%<Ó%>À!Ó%CØ#$¨¨qÑ1ÐØ"Ð#3·N±NÑ#CÐ"DÀAÐFˆCà�o‰o×!Ñ! #Ó+ˆØ×Ñ˜zÓ)Ø�U˜$˜t×/Ñ/Ð0°Ð2Ð2Øˆ
rˆ   c                ó^  • U R                   (       a  [        R                  (       a  gUR                  S:H  =(       aq    UR                  (       + =(       aY    U R
                  (       + =(       aA    [        R                  R                  R                  UR                  [        5       5      (       + $ )NFrN   )r<  r"   Ú!combo_kernel_per_subkernel_blocksrv  Úhas_zdimrQ  rL   rP  rH  r�  r  r4   )r  rå  s     r†   r  Ú#TritonKernel.needs_yz_grid_overflowá  sw   € ð ××¤F×$L×$LØà�N‰N˜aÑ÷ YØ—N‘NÔ"÷Yà×.Ñ.Ô.÷Yô —G‘G×$Ñ$×9Ñ9¸%¿+¹+Ä~ÓGWÓXÔXð		
rˆ   c                ó’   • U R                   (       a   U R                   UR                  5        S3   $ [        UR                  5          $ )Nr¾   )r  rõ   r2   )r  rÓ   s     r†   r  ÚTritonKernel.max_blockí  s;   € Ø××Ø×$Ñ$¨¯©«Ð'7°uÐ%=Ñ>Ð>Ü §¡£Ñ/Ð/rˆ   c                óÊ  • U R                   (       a:  [        R                  R                  R	                  UR
                  S5      (       a  gU R                  (       d  gU R                  (       a[  UR                  R                  5        S3U R                  ;   a0  U R                  UR                  R                  5        S3   S:X  a  gOKU R                  (       d:  [        R                  R                  R                  UR
                  S5      (       a  gUR                  (       a-  U R                  (       a  U R                  UR
                  5      nO?UR                  S:X  a  U R                  (       a  SnOU R!                  UR                  5      nUR                  (       a#  U R"                  (       a  X R%                  5       -  n[        R                  R                  R'                  UR
                  U5      (       ae  UR(                  S:g  =(       dO    UR*                  =(       d<    [        R                  R                  R-                  UR
                  [/        5       5      $ g)Nr  Fr¾   rN   TrN  )rE  rL   rP  rH  Ústatically_known_ltr  r;  r  rÓ   rõ   r<  rQ  r\  r  r  rV  r  rQ  r[  r�  rv  rï  r�  r4   )r  rá   r  s      r†   r'  ÚTritonKernel._has_constant_maskò  s£  € Ø× × ô �w‰w×Ñ×3Ñ3°D·J±JÀ×CÑCØà×!×!Øà×× D§K¡K×$5Ñ$5Ó$7Ð#8¸Ð!>À$×BSÑBSÓ!SØ× Ñ  D§K¡K×$5Ñ$5Ó$7Ð#8¸Ð!>Ñ?À1ÓDØð Eà×%×%Ü�w‰w×Ñ×7Ñ7¸¿
¹
ÀA×FÑFØð ×× ×!:×!:Ø×3Ñ3°D·J±JÓ?‰IØ�[‰[˜CÓ D§M§MØ‰IàŸ™ t§{¡{Ó3ˆIà×× ×!;×!;Ø!§O¡OÓ$5Ñ5ˆIô �7‰7×Ñ×8Ñ8¸¿¹ÀY×OÑOà—‘ Ñ"÷ WØ—=‘=÷Wä—7‘7×#Ñ#×8Ñ8¸¿¹Ä^ÓEUÓVðð rˆ   c                óf   • U R                   S   nUR                  S:X  d   eU R                  U5      $ )Nr   rN  )re  rÓ   r'  )r  Úxtrees     r†   r~  Ú TritonKernel._has_constant_xmask#  s5   € Ø× Ñ  Ñ#ˆØ�|‰|˜sÓ"Ð"Ð"Ø×&Ñ& uÓ-Ð-rˆ   c                ó¸   • U R                    H9  nU R                  U5      (       d  M  UR                  UR                   S35        M;     UR                  S5        g )Nr  r§   )re  r'  rX  rÓ   )r  rû   rá   s      r†   rÎ  ÚTritonKernel.filter_masks(  sL   € Ø×$Ô$ˆDØ×&Ñ& t×,Ó,Ø×!Ñ! T§[¡[ M°Ð"6Ö7ñ %ð
 	×Ñ˜&Õ!rˆ   c                óˆ   • [        [        R                  5      S U R                    Vs/ s H  n[        U   PM     sn$ s  snf rš   )ra  r¸   rð   rX  r   )r  rÝ   s     r†   Úget_reduction_prefixesÚ#TritonKernel.get_reduction_prefixes0  sG   € ô œ]×:Ñ:Ó;Ð<U¸d×>UÑ>UÑVó
âV�ô �tÔÙVñ
ð 	
ùò 
s   ©?c                óì  • U R                    Vs/ s H  o"R                  (       d  M  UPM     nnSR                  [        S U 5       5      5      nUR	                  SU R                  U5       35        U R                    Vs/ s H3  nUR                  (       d  M  [        R                  UR                     PM5     nn[        U5      nUR	                  SU R                  U5       35        gs  snf s  snf )zN
Generates code that flattens ND reduction numels, block sizes, etc. into 1D.
rz  c              3  ó>   #   • U  H  oR                    S 3v •  M     g7f)r  NrÍ  r„  s     r†   r  Ú8TritonKernel.codegen_reduction_numels.<locals>.<genexpr>=  s   é € Ð"UÂ_¸T§k¡k ]°%Õ#8Â_ùr„  z	rnumel = zRBLOCK: tl.constexpr = N)
re  r\  r   r!  rz  r6  r¸   ræ   rÝ   rD   )r  r&  rá   Úreduction_treesrÉ  Ú	rn_blocksrÐ  s          r†   rP  Ú%TritonKernel.codegen_reduction_numels7  sÓ   € ð
 -1×,<Ò,<ÓRÒ,< D×@QÕ@QŸ4Ñ,<ˆÐRØ—‘œFÑ"UÁ_Ó"UÓUÓVˆØ�‰˜	 $§*¡*¨VÓ"4Ð!5Ð6Ô7ð
 ×(Ò(ó
â(�Ø× Õ ó 1ŒM×%Ñ% d§i¡iÔ0Ù(ð 	ð 
ô
 ˜yÓ)ˆØ�‰Ð/°·
±
¸6Ó0BÐ/CÐDÕEùò Sùò

s   �C,§C,ÂC1Â!C1c                ó„   • U R                  5       nU Vs/ s H  n[        R                  " U U 340 UD6PM      sn$ s  snf )z;
Helper to initialize symbols like rn_numel, rn_base, etc.
)rü  r‹   rô   )r  r>  r   Úrn_prefixesrÓ   s        r†   rˆ  Ú#TritonKernel._get_reduction_symbolsI  sA   € ð ×1Ñ1Ó3ˆÙJUÓVÊ+À”—’ ˜x¨ xÐ0Ñ;°FÔ;É+ÑVÐVùÒVs   •%=c                óê   • U R                  5       nU R                  SSSS9n[        [        U5      S-
  5       Vs/ s H  n[	        X#S-   S 5      PM     sn[
        R                  " S5      /-   $ s  snf )z˜
Compute coefficients to convert ND reduction indices to linear indices.
For example:
  rindex = r0_index * r1_numel * ... * rn_numel + ... + rn_index.
r  Tr¿   rN   N)rü  rˆ  rc  r�   rD   r‹   r�   )r  r	  Ú	rn_numelsr…  s       r†   Ú_get_reduction_index_coeffsÚ(TritonKernel._get_reduction_index_coeffsP  s~   € ð ×1Ñ1Ó3ˆØ×/Ñ/°ÀÐPTÐ/ÐUˆ	ä;@ÄÀ[ÓAQÐTUÑAUÔ;Vó
Ú;V°CŒM˜)¨!¡G IÐ.Ö/Ñ;Vñ
ä�]Š]˜1ÓÐñð 	ùò 
s   »A0c                ó8   • U R                  5       n[        X!5      $ )z;
Compute linear reduction indices from N dimensional ones.
)r		  rC   )r  Ú
multi_indsÚcoeffss      r†   r‰  Ú'TritonKernel._flatten_reduction_indices]  s   € ð ×1Ñ1Ó3ˆÜ˜Ó,Ð,rˆ   c                ó  • U R                  SSSS9nU R                  SSSS9nU R                  U5      nUR                  SU R                  U5       35        U R                  U5      nUR                  SU R                  U5       35        g)zH
Generates code that converts ND reduction indices into linear indices.
rº   Tr»   rþ   z
roffset = z	rindex = N)rˆ  r‰  rz  r§  )r  r&  Ú
rn_offsetsÚrn_indsrs  Úrindexs         r†   rŠ  Ú&TritonKernel.codegen_reduction_indicesd  sš   € ð
 ×0Ñ0Ø˜d°ð 1ð 
ˆ
ð ×-Ñ-¨g¸tÐQUÐ-ÐVˆð ×1Ñ1°*Ó=ˆØ�‰˜
 4×#4Ñ#4°WÓ#=Ð">Ð?Ô@Ø×0Ñ0°Ó9ˆØ�‰˜	 $×"3Ñ"3°FÓ";Ð!<Ð=Õ>rˆ   c                ó>  • UR                   nUR                  (       a%  UR                  UR                   SU SU S35        OèUR                  cD  UR                  UR                   SU R                  U5       35        UR                  U S35        O—UR                  b  U SU R                  U5       3nOU R                  X S35      nU R                  (       d  UR                  5        S3OSnUR                  U SU R                  U5       S	U 3UR                   SU 3/5        U R                  U5      (       a!  UR                  U R                  U5      5        g US
:X  a  U R                  (       d%  UR                  U SUR                   SU S35        g g )Nr  z	offset + r$  z
offset = 0rº   r¾   r]  z	offset = rz  rN  rÕ  rÖ  r  )rÓ   rƒ  rh  rÈ   rv  r‡  rÖ   rå  rÚ  rõ   rÛ  rë  r'  Úcreate_constant_mask)r  rå  rV  rN  r%  Ú
block_sizes         r†   r†  Ú,TritonKernel.iteration_ranges_codegen_headert  s€  € ð �L‰LˆØ�=�=Ø�N‰N˜eŸj™j˜\¨¨Q¨C¨y¸¸¸4Ð@ÕAØ�^‰^Ñ#à�N‰N˜eŸj™j˜\¨¨T×-NÑ-NÈuÓ-UÐ,VÐWÔXØ�N‰N˜a˜S 
Ð+Õ,à×ÑÑ+Ø˜˜I d×&GÑ&GÈÓ&NÐ%OÐP‘à×8Ñ8¸À#ÀVÀÓM�ð ,0×+C×+C�1—7‘7“9�+˜UÑ#Èð ð �O‰Oà�c˜ 4×#@Ñ#@ÀÓ#GÐ"HÈÈJÈ<ÐXØ—z‘z�l # d VÐ,ðôð ×"Ñ" 5×)Ñ)Ø�N‰N˜4×4Ñ4°UÓ;Õ<Ø�s“(˜t×7×7à�N‰N˜a˜S ¨¯
©
 |°3°q°c¸Ð?Õ@ð  8rˆ   )rK  rM  rL  rŠ  rN  rE  rG  ry  rÆ   r  rU  r5  rI  rS  rY  r<  rÚ  r;  rC  rH  rA  rB  r@  r?  rx  rV  r   rO  )r   TNNF)
rW  zdict[str, sympy.Expr]r  zOptional[FixedTritonConfig]rI  zOptional[int]r<  r¥   r¦   r§   r%  ©rƒ  r¤  r¦   rÔ   r&  )rþ   rí   rš  z Optional[Union[str, tuple[str]]]rB  z!Optional[TMACompatibilityChecker])r°   )rÈ   rÔ   rÚ   rÔ   rD  z/Union[BlockPtrOptions, TensorDescriptorOptions]r¦   ztuple[str, str])r’   rí   rÞ  rí   r#  r¥   rõ   r¥   )rV  rV   )rÈ   rÔ   rq  rt  )rÈ   rÔ   rþ   rí   rš   )
rÈ   rÔ   rþ   rí   rª  rT   r�  rK   r¦   r§   r�  )r¡  úOptional[CSEVariable]©NN)r¿  rT   r¨  z.tuple[str, sympy.Expr, sympy.Expr, sympy.Expr]r©  rT   rª  r¤  r  r¥   r«  z Optional[tuple[str, sympy.Expr]]r¬  r	  r¦   rT   )r¦   ztuple[str, BlockShapeType])rª  rT   rƒ  r¤  r¦   rT   )
rƒ  r¤  rÕ  r¤  rƒ   rJ   rª  ú+Union[CSEVariable, tuple[CSEVariable, ...]]r¦   r	  )rƒ  r¤  r  )rÈ   rÔ   rþ   rí   rª  rT   )r¿  útuple[CSEVariable, ...]rƒ  útuple[torch.dtype, ...]r¦   rÔ   )rƒ  r	  r(  zUCallable[[tuple[CSEVariable, ...], tuple[CSEVariable, ...]], tuple[CSEVariable, ...]]r¿  r	  r¦   r	  )
rƒ  r	  r¿  r	  rÊ  r¥   rË  r¥   r¦   r	  )r¦   ry  )r  zOptional[float]r¦   rV   )r¦   z type[triton_heuristics.GridExpr]r  )rÈ   rÔ   r[  zOptional[IRNode]rÐ  r¥   )r¦   r–  )rå  rb   )rå  rc   r¦   rÔ   )rå  rc   rª  r	   r¦   rÔ   )rå  rc   r¦   r¥   )rÓ   rÔ   r¦   rî  )rá   rc   r¦   r¥   )rû   rú   r¦   r§   )r&  rV   r¦   r§   )r>  rÔ   r¦   zlist[sympy.Symbol]rµ  )r	  rO  r¦   rí   )rå  rc   rV  rV   r¦   r§   )hr›   r¨   r©   rª   r«   r3  rx  r¬   r  r6  rî  r   rm  rÈ  r7  r»  r8  r9  rš  r$  rc  r'  ri  rl  rp  rR  rT  rS  r�  r’  r•  r˜  rD  r  r  r'  r"  rG  rO  r0  rK  rR  r[  rr  rQ  r•  r›  rŽ  r¤  r´  r¼  r¿  rÄ  r7  r=  r
  rö  r  r  r[  r  rh  r†  r»  r£  rÐ  rô  rý  r  r$  r(  r­   r   r0  r7  rT  r®  r  r´  r’  r  rø  rÓ  rÙ  ra  rß  r‡  rå  rë  r  r  r'  r~  rÎ  r;   rü  rP  rˆ  r		  r‰  rŠ  r†  r®   r¥  r¦  s   @r†   r  r  Ç	  sw  ø‡ ñð &€IØ%Ó%Ø).€EÐ&Ó.Ø€OØ$;Ð!Ø3BÐÐ0ÓBàð "Ð#@ó ð KOÐ7¸ÓNð
 ØØ48Ø'+Ø %ð39à%ð39ð
 2ð39ð %ð39ð ð39ð 
÷39ð 39ðj ó:ó ð:ð6 ó
ó ð
ô
"ô
ò
#òJò*:ò0ô
ò

ð ó"ó ð"ð 8<ØØØØGKñb
àðb
ð 5ð	b
ð $Eõb
ðR ðf'àðf'ð ðf'ð Bð	f'ð 
õf'ôPXð:MàðMð ðMð ð	Mð
 ôMò4ð *€HØ8€Jàñ
ó ð
ð :?õ-ò<1ô  ð"
Øð
Ø:Hô
ôxðv SWð\Øð\Ø *ð\Ø3>ð\ØFOð\à	õ\ô|Iòô
ð& 48Ø04ð9àð9ð Cð9ð &ð	9ð
 $ð9ð ð9ð 1ð9ð .ð9ð 
õ9ôv.ô
9ð
Ø(ð
Ø1<ð
à	ô
ð(iàðið ðið &ð	ið
 ;ðið 
5ôiðV&ØDOô&ô"
ò0F
òP#ò0&ò!ð
à	ôðB>àð>ð ð>ð ô	>ð@>SØ1ð>SØ;Rð>Sà	ô>Sð@"à'ð"ð
ð"ð (ð"ð 
!ô"ðBD"à'ðD"ð (ðD"ð ð	D"ð
 ðD"ð 
!ôD"ôL$òm%ô^ô0wòr
ò
ð Ùˆtƒ_ó%"ó ó ð%"ðN ñ
ó ð
ð ñ-ó ð-ö^AðF
 ñó ðð  ñó ðò(;ôTCò(
-ð OSð)Øð)Ø/ð)ØGKõ)ô.,ô2ô	&ôJð*:Ø(ð:Ø14ð:à	ô:ôô.

ô0ô
/ôb.ô
"ð ó
ó ð
ôFô$Wð ó
ó ð
ô-ô?ð AØ(ðAØ0>ðAà	÷Aò Arˆ   r  c            
      óˆ  ^ • \ rS rSr% Sr\rS\S'   \" \	R                  \	R                  \	R                  \	R                  \	R                  \	R                  \	R                   \	R"                  /5      rSU 4S jjr\SS j5       rSS jrS rS	 rSSS
 jjr S   SS jjr        SS jr        SS jrS rSrU =r$ )ÚTritonSchedulingi”  z5Scheduling backend for Triton kernel code generation.z	type[Any]Úkernel_typec                óÄ   >• [         TU ]  U5        Ub  [        US5      (       d  g UR                   H+  n[	        U[
        [        45      (       d  M   [        Ul        M-     g )NrY  )	r™  rš  r³   rY  rŠ   r9   r7   Údebug_triton_codeÚdebug_device_str)r  rE  r[  rœ  s      €r†   rš  ÚTritonScheduling.__init__¥  sN   ø€ Ü‰Ñ˜Ô#ØÑ¤G¨I°w×$?Ñ$?ØØ—O”OˆDÜ˜$¤Ô0BÐ C×DÓDÜ(9�Ö%ò $rˆ   c                óä   • [         R                  R                  (       d  [         R                  R                  (       a'  [	        / U R
                  Q[        R                  P5      $ U R
                  $ rš   )r"   rµ   Úcooperative_reductionsÚforce_cooperative_reductionsr   Úbackend_featuresrQ   ÚREDUCE_TO_SINGLE_ELEMENT)rŸ   r
  s     r†   Úget_backend_featuresÚ%TritonScheduling.get_backend_features­  sR   € ô �M‰M×0×0Ü�}‰}×9×9äØP�#×&Ñ&ÐP¬×(OÑ(OÐPóð ð ×#Ñ#Ð#rˆ   c                ó  ^
• [         R                  R                  n[        X5      u  pEU(       a  UR	                  U5        [
        R                  (       aƒ  SSKJnJ	m
  [        U
4S jU 5       5      (       da  U Vs/ s H%  n[        Xv5      (       d  M  UR                  5       PM'     nnUR	                  UR                   SSR                  U5       35        U(       a  [        UU5      n	UR!                  X)5        g g s  snf )Nr   )r6   ÚForeachKernelSchedulerNodec              3  ó<   >#   • U  H  n[        UT5      v •  M     g 7frš   )rŠ   )rö   r¡  r-	  s     €r†   r  Ú3TritonScheduling.codegen_comment.<locals>.<genexpr>Ä  s   øé € ð ÚCP¸a”
˜1Ð8×9Ð9Â=ùs   ƒz Fused node name list: r   )rL   rP  rÆ  r?   Úmake_commentr"   Údebug_fusionÚtorch._inductor.schedulerr6   r-	  rz   rŠ   Úget_nameÚcommentr   r+   Úwrite_provenance_debug_handle)r  Únode_schedulerc  rÑ  ÚoriginsÚ_detailed_originsr6   r¡  Ú
node_namesÚdebug_handler-	  s             @r†   Úcodegen_commentÚ TritonScheduling.codegen_comment¸  sé   ø€ Ü—'‘'×&Ñ&ˆÜ%8¸Ó%PÑ"ˆÞØ× Ñ  Ô)ä××÷ô
 ô ÙCPó÷ ñ ñ +óâ*˜Ü! !×7ó !�A—J‘J–LÙ*ð ð ð
 ×$Ñ$Ø—‘Ð'Ð'>¸t¿y¹yÈÓ?TÐ>UÐVôö ÜBØØóˆLð ×1Ñ1°+ÕLð ùòs   Á<C?ÂC?c	                óì  • UR                  UUUUUU5      (       a  g[        5       n	[        R                  5       (       a  [        R	                  XS5        U	R                  SU< S35        U	R                  USS9  [        R                  R                  5       n
U	R                  SU
R                   S35        SU 3nU" Xa5      u  pÍUS	U-   S	-   U-   -  nUR                  XIR                  5       U5        g)
zDEmit kernel to wrapper, with support for external template handlers.Nzasync_compile.triton(z, '''Tr^  z''', device_str='z')z# kernel path: r  )Úemit_kernel_overriderV   Úasync_compileÚuse_process_poolrµ   rh  rz  rL   rP  r
  rA  Údefine_kernelr‚  )r  rÑ  rÅ   Úsrc_coderc  Ú	subs_namer6	  Úkernel_pathr?   Úcompile_wrapperr  Úmetadata_commentr7	  Údetailed_originss                 r†   Ú_emit_kernel_to_wrapperÚ(TritonScheduling._emit_kernel_to_wrapperÙ  sÿ   € ð ×&Ñ&ØØØØØØ÷
ñ 
ð ä(Ó*ˆä×)Ñ)×+Ñ+ô × Ñ  Ô5à×!Ñ!Ð$9¸)¹ÀeÐ"LÔMà×Ñ˜x¨tÐÑ4ÜŸ™×<Ñ<Ó>ˆØ×!Ñ!Ð$5°n×6IÑ6IÐ5JÈ"Ð"MÔNà,¨[¨MÐ:ÐÙ$7¸Ó$OÑ!ˆØ˜D 7™N¨TÑ1Ð4DÑDÑDÐØ×Ñ˜k×+CÑ+CÓ+EÐGWÕXrˆ   c           
     óò  • [         R                  R                  nXR                  ;   a  UR                  U   nU$ [        R
                  R                  (       a$  [        U[        R
                  R                  5      OSn[        U5      S S nSR                  SXvUR                  5       /5      n[        R                  R                  (       a  [        R                  R                   SU 3nXTR                  U'   [        R
                  R                  (       a  UOSnUR                  [        [         R"                  5      U5      nUR                  [        [         R$                  5      U5      nUR                  SS5      n['        [)        UR+                  5       5      S5      u  pšnU R-                  UUUUUUU[.        5        [0        R2                  " S	5      (       a  [0        R4                  " X[U5        U$ )
Nr°   r   rW  rµ   Útriton_z#pragma CMTÚ#ÚpyÚkernel_metadata)rL   rP  rÆ  Úsrc_to_kernelr"   rµ   Údescriptive_namesr>   rM   r   Únext_kernel_suffixÚaot_inductorÚmodel_name_for_generated_filesÚunique_kernel_namesr  rÔ   rA   r�  r  r(   r'   r_  rH	  r?   r$   Úis_metric_table_enabledÚlog_kernel_metadata)r  rB	  r6	  rÅ   rÑ  rc  Ú
fused_nameÚkernel_categoryrC	  Ú	_basenamerW  rD	  s               r†   rA	  ÚTritonScheduling.define_kernel  s¬  € Ü—'‘'×&Ñ&ˆØ×,Ñ,Ó,Ø!×/Ñ/°Ñ9ˆKðb Ðô[ —=‘=×2×2ô & m´V·]±]×5TÑ5TÔUàð ô
 AÀÓJÈ2ÈAÐNˆOØŸ(™(Ø˜?¸×8RÑ8RÓ8TÐUóˆKô ×"Ñ"×A×Aô "(×!4Ñ!4×!SÑ!SÐ TÐTUÐVaÐUbÐc�ð /:×!Ñ! (Ñ+Ü'-§}¡}×'H×'H™ÈiˆIð
  ×'Ñ'¬¬K×,HÑ,HÓ(IÈ;ÓWˆHØ×'Ñ'¬¬K×,CÑ,CÓ(DÀiÓPˆHð  ×'Ñ'¨°sÓ;ˆHä(0´¸8¿>¹>Ó;KÓ1LÈdÓ(SÑ%ˆI˜+à×(Ñ(ØØØØØØØÜ#ô	ô ×.Ò.Ð/@×AÑAÜ×+Ò+¨KÀhÔOàÐrˆ   c                óŽ   • U R                  USS9n[        R                  " U5      nU R                  XB[	        S U 5       5      S9$ )zk
Benchmark fused list of nodes and return the execution time
in milliseconds on randomly generated inputs.
T)rt  c              3  ó@   #   • U  H  oR                  5       v •  M     g 7frš   ©r3	  ©rö   r¡  s     r†   r  Ú9TritonScheduling.benchmark_fused_nodes.<locals>.<genexpr>A  s   é € Ð:WÒQVÈA¿:¹:¿<¸<ÒQVùrÈ  )r9	  )Úgenerate_kernel_code_from_nodesr)   rQ  Úbenchmark_codegened_moduler   )r  rY  Ún_spills_thresholdrB	  rþ  s        r†   Úbenchmark_fused_nodesÚ&TritonScheduling.benchmark_fused_nodes9  sS   € ð
 ×7Ñ7¸ÐPTÐ7ÐUˆÜ×Ò˜xÓ(ˆØ×.Ñ.Ø´
Ñ:WÑQVÓ:WÓ0Wð /ð 
ð 	
rˆ   c           	     óf  ^^
^^^^• [        [        R                  R                  5      n[	        5          UR                  [        R                  R                  5       5         SmU4S jmUU4S jnU4S jnUb  UO[        S/5      n[        R                  SUTR                  5        U" 5       mTb   TTR                  4sSSS5        sSSS5        $ TR                  5       m
TR                  mTR                  m T" TR                  " T
6 S   5        TR(                  n[+        U5      S
:X  d   eUS   R,                  U:”  a  ['        S	5      mOx[        R                  R                  5       n	[.        R0                  " U
UU4S jU	S9m[+        TR2                  5      S:”  a&  T[.        R0                  " U
U4S j[5        U	5      S9-
  m[        R                  SUT5        U" 5         TTR                  4sSSS5        sSSS5        $ ! [         as  n[         R"                  R$                  (       a  e [        R                  SUU5        ['        S	5      mU" 5         TTR                  4s SnAsSSS5        sSSS5        $ SnAff = f! , (       d  f       O= f SSS5        g! , (       d  f       g= f)z$Benchmark an already compiled moduleNc                 ó€   >• T R                   c   e[        R                  R                  T R                   5      S   S-   $ ©Nr   z.kernel_perf©Ú__file__ÚosÚpathÚsplitext©rþ  s   €r†   Úcache_file_pathÚDTritonScheduling.benchmark_codegened_module.<locals>.cache_file_pathO  s6   ø€ Ø—|‘|Ñ/Ð/Ð/Ü—w‘w×'Ñ'¨¯©Ó5°aÑ8¸>ÑIÐIrˆ   c                 ó>   >• T" 5       n [        U [        T5      5        g rš   ©r*   rÔ   )rk	  rn	  Úmss    €€r†   Ústore_cacheÚ@TritonScheduling.benchmark_codegened_module.<locals>.store_cacheS  s   ø€ Ù&Ó(�Ü˜T¤3 r£7Õ+rˆ   c                 óÚ   >• T" 5       n [         R                  R                  U 5      (       a.  [        U 5       n[	        UR                  5       5      sS S S 5        $ g ! , (       d  f       g = frš   )rj	  rk	  ÚexistsÚopenr,  Úread©rk	  Úfdrn	  s     €r†   Ú
load_cacheÚ?TritonScheduling.benchmark_codegened_module.<locals>.load_cacheW  sJ   ø€ Ù&Ó(�Ü—7‘7—>‘> $×'Ñ'Ü˜dœ rÜ$ R§W¡W£YÓ/÷ $™à÷ $œàús   ¸AÁ
A*rO  ú%kernel src code for %s written to: %sr   z*Exception (%s) in compiling fused nodes %sÚinfrN   c                 ó4   >• T" TR                   " T 6 S   5      $ rx  ©Ú
clone_args©rÂ  ÚcallÚwrapped_jit_functions   €€€r†   rÒ  Ú=TritonScheduling.benchmark_codegened_module.<locals>.<lambda>ˆ  ó   ø€ ™DÐ!5×!@Ò!@À$Ð!GÈÑ!JÔKrˆ   ©r
  c                 ó"   >• TR                   " T 6 $ rš   r€	  ©rÂ  r„	  s   €€r†   rÒ  r…	  ‘  s   ø€ Ð 4× ?Ò ?ÀÑ Frˆ   z+The fused kernel for %s took %.3f ms to run)r   rL   rP  r—  r   r
  r
  r   rø  r  ri	  Úget_argsrƒ	  rK	  r�	  Ú	Exceptionr"   rµ   Ú.disallow_failing_autotune_kernels_TESTING_ONLYr,  Ú	launchersr�   Ún_spillsr.   Ú	benchmarkrd  rÔ   )r  rþ  rb	  r9	  Údevice_interfacers	  r{	  r"  r�	  r
  rÂ  rn	  rƒ	  rr	  r„	  s    `        @@@@@r†   ra	  Ú+TritonScheduling.benchmark_codegened_moduleD  sK  ý€ ô 4´A·G±G×4GÑ4GÓHÐäÕ Ø×#Ñ#¤A§G¡G×$GÑ$GÓ$IÕJàˆBõJö,õð )Ñ4‘
¼*ÀiÀ[Ó:Qð ô �I‰IØ7ØØ—‘ôñ
 “ˆBØ‰~Ø˜3Ÿ<™<Ð'÷; KÐJ÷ !Ñ ð@ —<‘<“>ˆDØ—8‘8ˆDØ#&§;¡;Ð ð(ÙÐ)×4Ò4°dÐ;¸AÑ>Ô?ð -×6Ñ6ˆIÜ�y“> QÓ&Ð&Ð&ð ˜‰|×$Ñ$Ð'9Ó9ä˜5“\‘äŸ™×<Ñ<Ó>�ô !×*Ò*ÞKØ!ñ�ô Ð+×=Ñ=Ó>ÀÓBØœk×3Ò3ÝFÜ" 6›{ññ �Bô
 �I‰IØ=ØØôñ
 ŒMØ�s—|‘|Ð#÷a KÐJ÷ !Ñ øôL ó 
(Ü—=‘=×O×OØÜ—	‘	Ø@ØØôô
 ˜5“\�Ù”Ø˜3Ÿ<™<Ð'Ô'÷_ KÐJ÷ !Ñ ûðL
(ú÷K KÕJúÐJ÷ !× Ö úsh   ´.J"Á"AJÃ	J"Ã(JÃ=HÄCJÇ4	J"È
JÈAI?É'JÉ(JÉ,	J"É?JÊJÊ
J	ÊJ"Ê"
J0c                óú  • UR                  S5      nU=(       a     [        S UR                  5        5       5      nU R                  nU(       a  SSKJn  UnU(       a  SUS'   UR                  S5      (       a
  SUS	'   SUS'   [        R                  UR                  5      (       d  UR                  S	5      (       a   eSUS	'   [        R                  R                  XaX#5      nU" U0 UD6nU R                  X‚U5      $ )
Nr»  c              3  ó@   #   • U  H  oR                  5       v •  M     g 7frš   )Úis_split_scan)rö   r[  s     r†   r  Ú9TritonScheduling.create_kernel_choices.<locals>.<genexpr>¤  s   é € ð (
Ú-N T×Ñ× Ð Ò-NùrÈ  rN   )ÚTritonSplitScanKernelFÚoverride_cooperative_reductionr£  TÚoverride_persistent_reduction)Úcontains_oprz   Úscheduler_nodesr 	  Útriton_split_scanr–	  r  r´  rŽ  rè  rL   ro  Útriton_kernel_kwargsÚadd_multi_kernel_choices)	r  Úkernel_featuresÚkernel_argsÚkernel_kwargsÚis_scanr”	  r 	  r–	  rÅ   s	            r†   Úcreate_kernel_choicesÚ&TritonScheduling.create_kernel_choices�  s  € ð "×-Ñ-¨fÓ5ˆØ÷ 
¤Cñ (
Ø-<×-LÑ-LÔ-Nó(
ó %
ˆð +/×*:Ñ*:ˆÞÝ@à/ˆKæà>CˆMÐ:Ñ;ð ×&Ñ& v×.Ñ.Ø=AˆMÐ9Ñ:Ø>CˆMÐ:Ñ;ä×1Ñ1°/×2QÑ2Q×RÑRà$×(Ñ(Ð)H×IÑIÐIÐIØ=BˆMÐ9Ñ:äŸ	™	×6Ñ6Ø¨+ó
ˆñ ˜kÐ;¨]Ñ;ˆØ×,Ñ,¨VÀ-ÓPÐPrˆ   c           	     ó$  • U/n[         R                  R                  (       d  U$ UR                  =(       a    UR	                  S5      (       + nUR
                  =(       a    UR	                  S5      (       + nU(       a%  UR                  U R                  " U0 UDSS0D65        U(       aª  UR                  R                  n[        R                  R                  R                  US5      (       ae  UR                  U R                  " U0 UDSS0D6=n5        U(       a7  UR                  (       a&  UR                  U R                  " U0 UDSSS.D65        [        U5      S:”  a-  USS   H  n	UR                  U	l        M     UR!                  S S9  U$ )	Nr˜	  r—	  Fi   )r—	  r˜	  rN   c                ó   • U R                   $ rš   )r  )r*  s    r†   rÒ  Ú;TritonScheduling.add_multi_kernel_choices.<locals>.<lambda>ô  s
   €  q×'=Ò'=rˆ   rž  )r"   rµ   Úmulti_kernelr  rè  rQ  rØ  r 	  r[  rŽ  rL   rP  rH  r�  r�   Úmust_keep_buffersr£  )
r  rÅ   rŸ	  r 	  ÚkernelsÚoptional_persistentÚoptional_cooperativerÉ  rf  Úkernel2s
             r†   r�	  Ú)TritonScheduling.add_multi_kernel_choicesÂ  s‹  € ð (. hˆÜ�}‰}×)×)ØˆNà$×9Ñ9÷ 
À-×BSÑBSØ+óC
ô ?
Ðð  &×;Ñ;÷  
ÀM×DUÑDUØ,óE
ô A
Ðö Ø�N‰NØ× Ò Ø ðà#ñð 38òôö  Ø—_‘_×4Ñ4ˆFä�w‰w×Ñ×4Ñ4°V¸U×CÑCØ—‘Ø!×-Ò-Ø$ðà'ñð 8=òð �Eôö '¨5×+E×+EØ—N‘NØ×(Ò(Ø(ðà+ðð <AØ:?ó	ôô ˆw‹<˜!ÓØ" 1 2›;�à,2×,DÑ,D�Ö)ñ 'ð �L‰LÑ=ˆLÑ>Øˆrˆ   c                ó¤  ^^^^^^^• U4S jmU4S jnUUU4S jnS/ peSn[         R                  R                  n[        U5      [         R                  l        [         R                  R                  n	[        U	5      [         R                  l        [
        R                  S:„  n
[
        R                  S:„  nU R                  USU
USS9nU GHH  u  pÞnU Vs/ s H  nUR                  5       PM     nnU VVs/ s H  nU  H  nUR                  5       PM     M     nnn[        U5      S:X  a'  X/S      u  nnUU-  nUS-  nUR                  U5        M‰  Uc   eUR                  [        [        R                   5      S	5      n["        R$                  " U5      m[&        R)                  S
UTR*                  5        U" 5       u  mmTb(  UT-  nUT-  nUR                  TR*                  5        GM#  TR-                  5       mTR.                  mTR0                  mT" TR2                  " T6 S   5        TR4                  n[        U5      S:X  d   eUS   R6                  S:”  a  [9        S5      =mmOS[         R                  R;                  5       n[<        R>                  " UUU4S jUS9m[<        R>                  " UU4S jUS9m[&        R)                  S[        S U 5       5      TT5        U" 5         UT-  nUT-  nUR                  TR*                  5        GMK     U[         R                  l        U	[         R                  l        XWU4$ s  snf s  snnf )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                 ó€   >• T R                   c   e[        R                  R                  T R                   5      S   S-   $ rg	  rh	  rm	  s   €r†   rn	  Ú@TritonScheduling.benchmark_combo_kernel.<locals>.cache_file_path  s6   ø€ Ø—<‘<Ñ+Ð+Ð+Ü—7‘7×#Ñ# C§L¡LÓ1°!Ñ4°~ÑEÐErˆ   c                 ó  >• T" 5       n [         R                  R                  U 5      (       aC  [        U 5       n[	        S UR                  5       R                  5        5       5      sS S S 5        $ g! , (       d  f       g= f)Nc              3  ó8   #   • U  H  n[        U5      v •  M     g 7frš   )r,  )rö   r"  s     r†   r  ÚNTritonScheduling.benchmark_combo_kernel.<locals>.load_cache.<locals>.<genexpr>  s   é € Ð EÒ3D¨a¤ q§ Ò3Dùrh  r	  )rj	  rk	  rv	  rw	  r×   rx	  Úsplitry	  s     €r†   r{	  Ú;TritonScheduling.benchmark_combo_kernel.<locals>.load_cache	  sW   ø€ Ù"Ó$ˆDÜ�w‰w�~‰~˜d×#Ñ#Ü˜$”Z 2Ü Ñ E°2·7±7³9·?±?Ô3DÓ EÓE÷  ‘Zà÷  ”Zàús   ¸.A1Á1
A?c                 ó\   >• T" 5       n [        U [        T5      S-   [        T5      -   5        g )Nr?  rq	  )rk	  rn	  rr	  Úms_clones    €€€r†   rs	  Ú<TritonScheduling.benchmark_combo_kernel.<locals>.store_cache  s&   ø€ Ù"Ó$ˆDÜ˜œs 2›w¨™}¬s°8«}Ñ<Õ=rˆ   r   g        T)Úsubkernel_nodesÚcustom_part_algorithmÚenable_autotuneÚmixed_sizesÚonly_gen_src_coderN   rK	  r}	  r~	  c                 ó4   >• T" TR                   " T 6 S   5      $ rx  r€	  r‚	  s   €€€r†   rÒ  Ú9TritonScheduling.benchmark_combo_kernel.<locals>.<lambda>S  r†	  rˆ   r‡	  c                 ó(   >• TR                   " T 6 S   $ rx  r€	  r‰	  s   €€r†   rÒ  r¿	  W  s   ø€ Ð0×;Ò;¸TÐBÀ1ÒErˆ   zDThe fused kernel for %s took %.3f ms to run, %.3f ms to clone inputsc              3  ó@   #   • U  H  oR                  5       v •  M     g 7frš   r]	  r^	  s     r†   r  Ú:TritonScheduling.benchmark_combo_kernel.<locals>.<genexpr>]  s   é € Ð<²¨AŸ:™:Ÿ<˜<²ùrÈ  ) rL   rP  rx  r   Úinplaced_to_remover"   Úcombo_kernels_autotuneÚcombo_kernel_allow_mixed_sizesÚgenerate_combo_kernel_codeÚ	get_nodesr3	  r�   rØ  r  rÔ   rA   r  r)   rQ  rø  r  ri	  rŠ	  rƒ	  rK	  r�	  r�	  rŽ	  r,  r
  r.   r�	  ) r  Ú	node_listÚnode_benchmark_resultsr{	  rs	  Útotal_msÚ	file_listÚtotal_clone_msÚremoved_buffers_origÚinplaced_to_remove_origr»	  r¼	  Úkernel_code_listrB	  rÅ   Ú
node_groupr[  Úfused_node_listsrY  r¡  ÚnamesÚnode_msrk	  r�	  r
  rÂ  rn	  rƒ	  rþ  rr	  r·	  r„	  s                             @@@@@@@r†   Úbenchmark_combo_kernelÚ'TritonScheduling.benchmark_combo_kernel÷  s  þ€ õ	Fõ	 ÷	>ð   �)Ø #ˆÜ Ÿw™w×6Ñ6ÐÜ",Ð-AÓ"BŒ�‰ÔÜ"#§'¡'×"<Ñ"<ÐÜ%/Ð0GÓ%HŒ�‰Ô"Ü ×7Ñ7¸!Ñ;ˆÜ×;Ñ;¸aÑ?ˆØ×:Ñ:Ø%Ø"&Ø+Ø#Ø"ð ;ð 
Ðô -=Ñ(ˆH˜jÙ=GÓHºZ°T §¡Ö 0¹ZÐÐHÙ/?ÔOÒ/? eÌÀA�Q—Z‘Z–\É‘\Ñ/?ˆEÑOä�:‹ !Ó#à 6À!±}Ñ E‘�˜à˜GÑ#�Ø !Ñ#�Ø× Ñ  Ô&ÙàÑ'Ð'Ð'Ø×'Ñ'¬¬K×,CÑ,CÓ(DÀiÓPˆHÜ×"Ò" 8Ó,ˆCä�I‰IØ7ØØ—‘ôñ
 &›<‰LˆB�Ø‰~Ø˜B‘�Ø (Ñ*�Ø× Ñ  §¡Ô.Úà—<‘<“>ˆDØ—8‘8ˆDØ#&§;¡;Ð ñ Ð%×0Ò0°$Ð7¸Ñ:Ô;à,×6Ñ6ˆIÜ�y“> QÓ&Ð&Ð&Ø˜‰|×$Ñ$ qÓ(ä % e£Ð,�‘XäŸ™×<Ñ<Ó>�ô !×*Ò*ÞKØ!ñ�ô '×0Ò0ÝEØ!ñ�ô
 �I‰IØVÜÑ<±Ó<Ó<ØØô	ñ ŒMØ˜‰NˆHØ˜hÑ&ˆNØ×Ñ˜SŸ\™\×*ñ -=ð@ #7Œ�‰ÔØ%<Œ�‰Ô"Ø¨Ð2Ð2ùòC  IùÛOs   ÃMÃ8$Mr‚   )rE  zOptional[Scheduler]r¦   r§   )r
  ztorch.devicerš   )rx   )r¦   útuple[float, str])rx   N)r9	  zOptional[OrderedSet[str]]r¦   rÖ	  )rž	  rs   rŸ	  ú	list[Any]r 	  rt  r¦   úlist[TritonKernel])rÅ   r  rŸ	  r×	  r 	  rt  r¦   rØ	  ) r›   r¨   r©   rª   r«   r  r 	  r¬   r   rQ   ÚFOREACHÚ	BUCKETIZEÚINPLACE_BUFFERSÚMASKED_SCATTER_WITH_INDEXÚSCANÚSORTÚTRITON_TEMPLATESÚTUPLE_REDUCTIONr(	  rš  r­   r*	  r;	  rH	  rA	  rc	  ra	  r¢	  r�	  rÔ	  r®   r¥  r¦  s   @r†   r	  r	  ”  s#  ø‡ Ù?à)€K�Ó)Ù!à×"Ñ"Ø×$Ñ$Ø×*Ñ*Ø×4Ñ4Ø×ÑØ×ÑØ×+Ñ+Ø×*Ñ*ð		
óÐ÷:ð ó$ó ð$ôMòB(YòT4öl	
ð RVðW$Ø5NðW$à	õW$ðr#Qà+ð#Qð ð#Qð &ð	#Qð
 
ô#QðJ3àð3ð ð3ð &ð	3ð
 
ô3÷jp3ð p3rˆ   r	  c                ó4  • / nU R                  5       nUb!  [        U[        R                  5      (       d   eU(       a1  UR                  c$  UR                  U R                  5        S35        U$ SSKJn  U R                  5       nUc   eU R                  R                  U5      n[        U[        U45      (       d   S[        U5       35       e[        R                  R!                  U5         ["        R$                  nUR'                  U R)                  5       5      R+                  5       nU["        l        S S S 5        UR                  U R                  5        S35        UR                  [,        R.                  " WS5      5        U$ ! , (       d  f       NX= f)Nz" Unfinalized multi template bufferr   )ÚCUDACombinedSchedulingz]Scheduling backend should be SIMD or CUDACombined when generating debug Triton strings, got: z Triton code:z    )Úget_template_noderŠ   r#   ÚMultiTemplateBufferÚmake_kernel_renderrØ  r3	  Ú0torch._inductor.codegen.cuda_combined_schedulingrâ	  r  rE  Úget_backendrf   rA  rL   rP  Úset_current_devicer$   Úgenerated_kernel_countr`	  rÇ	  r_  r!  rž  )r[  ÚlinesÚmulti_templaterâ	  r
  ÚbackendÚold_generated_kernel_countÚtriton_codes           r†   r"	  r"	  j  s_  € Ø€EØ×+Ñ+Ó-€NØÑ!¤Z°Ä×@VÑ@V×%WÑ%WÐWÐWÞ˜.×;Ñ;ÑCØ�‰˜Ÿ™›Ð(Ð(JÐKÔLð2 €Lõ/	
ð —‘Ó"ˆØÑ!Ð!Ð!Ø—.‘.×,Ñ,¨VÓ4ˆÜ˜'¤NÐ4JÐ#K×LÑLð 	
ØkÔlpÐqxÓlyÐkzÐ{ó	
ÐLô �W‰W×'Ñ'¨Õ/ô *1×)GÑ)GÐ&Ø!×AÑAØ—‘Ó óç‰e‹gð ð .HŒGÔ*÷ 0ð 	�‰˜Ÿ™›Ð(¨Ð6Ô7Ø�‰”X—_’_ [°&Ó9Ô:Ø€L÷ 0Õ/ús   Ã.A	F	Æ	
F)r’   rß  r&  )rª  rÔ   r±  r/  rÙ  r/  r¦   rÔ   r	  )rƒ  r¤  r¦   r¤  )rƒ  r¤  r¦   r¥   )rÚ   zUnion[CSEVariable, Any]r¦   r¥   )r¦   ro   rÅ  )r¡   r¥   r¦   zCallable[[_T], _T])r[  r6   r¦   ry  )ÐÚ
__future__r   rØ  r‹  r©  r6  r½  Úloggingr!  rë  rj	  r!  Úabcr   Úcollections.abcr   r   r   r   Útypingr	   r
   r   r   r   r   r‹   Úsympy.printing.precedencer   rœ   Útorch._loggingÚtorch.utils._pytreer%   Ú_pytreer  Útorch._dynamo.device_interfacer   Útorch._dynamo.utilsr   r   Útorch._prims_commonr   Útorch.utils._ordered_setr   Útorch.utils._sympy.functionsr   r   r   Útorch.utils._tritonr   r   r   Úutils._sympy.symbolr   r   r   r   Úutils._sympy.value_rangesr    r°   r"   r#   r$   r?	  r&   Ú	codecacher'   r(   r)   r*   r  r+   Úops_handlerr,   Úruntimer-   Úruntime.benchmarkingr.   Úruntime.hintsr/   r0   r1   r2   r3   Úruntime.runtime_utilsr4   r5   rE  r6   r7   r8   r9   Úshape_propagationr:   r;   r<   r=   r>   r?   r@   rA   rB   rC   rD   rE   rF   rG   rH   ÚvirtualizedrI   r  rJ   rK   rL   Úwrapper_benchmarkrM   Úblock_analysisrO   ÚcommonrP   rQ   rR   rS   rT   rU   rV   rW   rX   rY   rZ   r[   r\   r]   r^   r_   Úsimdr`   ra   rb   rc   rd   re   rf   Útriton_utilsrg   rh   ri   rj   rk   rÑ  rl   Útypesrm   r¨  ro   Ú%torch.fx.experimental.symbolic_shapesrp   rq   rr   Úsimd_kernel_featuresrs   rt   Ú	getLoggerr›   rø  Ú_loggingÚgetArtifactLoggerÚperf_hint_logÚschedule_logÚ
fusion_logr‡  r‡   r“   r•   r¶   r¸   rÒ  rø   r)  r»  rÈ  r¥  rà  r2  r  r„  r‡  r�  r�  r’  r”  r–  r©  rÎ  rÐ  Ú_initialize_pointwise_overridesr3  rw  r*  rÔ  rñ  rÔ   r×   rú  r   r  r	  r"	  r‚   rˆ   r†   Ú<module>r
     s»  ðå "ã Û Û Û Û Û Û Û Û 	Û Ý ß 8Ñ 8Ý ß E× Eã Ý 0ã Û ß $Ð $Ý Cß <Ý 0Ý /ß KÑ K÷ñ ÷ YÓ XÝ 4ß )Ó )Ý (ß FÓ FÝ ;Ý (Ý 'Ý .÷õ ÷ Dß WÓ WÝ 5÷÷ ÷ ÷ ÷  CÓ BÝ BÝ /÷÷ ÷ ÷ ó ÷$÷ ñ ÷õ õ %ö Ý åLÝ>åÝ&Ý8á	�‹€Bà×Ó˜Ó!€Ø—’×0Ò0°¸<ÓH€Ø�~Š~×/Ò/°¸*ÓE€Ø�^Š^×-Ò-¨h¸ÓA€
Ù“€ð #$Ð ò,ô	÷6ñ 6ñ  ˆ4ƒóó ð÷$Q,ñ Q,ðh ×Ò÷
ð 
ó ð
ðB ×Ò÷~ð ~ó ð~ðB
 ×Òô?Ð4ó ?ó ð?ð8 ×Òô=Ð,ó =ó ð=ð@+Øð+Ø/ð+Ø<Pð+àô+ô>XQ�Mô XQñv 	‹×Ò€ô3ô
ô&ô8ô
;ôPô˜ô ô8(ö.ôbu
&�kô u
&ðp × /Ò /°Ô 9ôH"˜Oô H"÷V$+ñ $+ðN ×Ò÷N
ð N
ó ðN
÷b!ñ !ðH ×Ò÷#ð #ó ð#ô
�Ð% u¨S°%¸¸S¸±/Ð-AÑ'BÐBÑCô 
ð ×Ò÷@ð @ó ð@ôFJ7A�:Ð/Ñ0ô J7AôZnS3�~ô S3õlrˆ   