Ë
    g^(h�§ ã                  óÚ  — d dl mZ d dlZd dlZd dlZd dlZd dlZd dlZd dlZd dl	Z	d dl
Z
d dlmZmZ d dlmZ d dlmZmZmZmZmZmZ d dlZd dlmZ d dlZd dlZd dlmc mZ d dlm Z  d dl!m"Z"m#Z# d d	l$m%Z% d d
l&m'Z' d dl(m)Z)m*Z*m+Z+ d dl,m-Z- ddl.m/Z/m0Z0m1Z1m2Z2 ddl3m4Z4 ddl5m6Z6m7Z7m8Z8 ddl9m:Z: ddl;m<Z<m=Z=m>Z> ddl?m@Z@ ddlAmBZB ddlCmDZD ddlEmFZFmGZGmHZHmIZI ddlJmKZKmLZL ddlMmNZNmOZOmPZPmQZQ ddlmRZRmSZSmTZTmUZUmVZVmWZWmXZXmYZYmZZZm[Z[m\Z\m]Z]m^Z^m_Z_ ddl`maZbmcZcmdZdmeZe ddlfmgZg ddlhmiZi ddljmkZkmlZlmmZmmnZnmoZompZpmqZqmrZrmsZsmtZtmuZumvZvmwZwmxZxmyZy dd lzm{Z{m|Z|m}Z}m~Z~mZm€Z€ dd!l�m‚Z‚mƒZƒm„Z„m…Z…m†Z† dd"l‡mˆZˆ er&d d#l‰mŠZŠ d d$lm‹Z‹ d d%lŒm�Z� dd&l7mŽZŽ dd'l�m�Z�  e‹d(«      Z‘ e�j$                  e“«      Z”e�j*                  �j-                  e“d)«      Z—e�j*                  �j-                  e“d*«      Z˜e�j*                  �j-                  e“d+«      Z™ e:«       Z9 G d,„ d-«      Zš ed«      dXd.„«       Z› ed«      dXd/„«       Zœ G d0„ d1«      Z�e�j<                   G d2„ d3«      «       ZŸe�j<                   G d4„ d5«      «       Z 	 	 	 	 	 	 	 	 dYd6„Z¡ G d7„ d8et«      Z¢ e¢«       �jF                  Z¤dZd9„Z¥d[d:„Z¦dZd;„Z§d\d<„Z¨dZd=„Z©d]d>„Zªd^d?„Z« G d@„ dAeo«      Z¬d_dB„Z­d`dadC„Z® G dD„ dEes«      Z¯e¯�ja                  dF«        G dG„ dHe¯«      Z± G dI„ dJ«      Z²e�j<                   G dK„ dL«      «       Z³ G dM„ dN«      Z´e�j<                   G dO„ dP«      «       Zµ G dQ„ dRene¬ee¶e·e¶e¶f   f   f   «      Z¸ G dS„ dTee¬   «      Z¹ G dU„ dVe€«      ZºdbdW„Z»y)cé    )ÚannotationsN)ÚIterableÚSequence)Ú	lru_cache)ÚAnyÚCallableÚcastÚOptionalÚTYPE_CHECKINGÚUnion)Ú
PRECEDENCE)Úget_interface_for_device)ÚidentityÚpreserve_rng_state)Úis_integer_dtype)Ú
OrderedSet)ÚCeilDivÚFloorDivÚModularIndexing)Úhas_triton_packageé   )Úfree_symbol_is_typeÚ
prefix_strÚsymbol_is_typeÚSymT)ÚValueRangesé   )ÚconfigÚirÚmetrics)ÚAsyncCompile)Ú	code_hashÚget_pathÚPyCodeCache)ÚDefaultHandler)Útriton_heuristics)Úbenchmarker)ÚAutotuneHintÚDevicePropertiesÚTRITON_MAX_BLOCKÚTRITON_MAX_RSPLIT)Úget_max_y_gridÚnext_power_of_2)ÚBaseSchedulerNodeÚFusedSchedulerNodeÚ	SchedulerÚSchedulerNode)Ú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ÚOpOverridesÚPythonPrinterÚ
RemovedArgÚSizeArgÚ	TensorArgÚWorkspaceArgÚWorkspaceZeroMode)Úconstant_reprÚIterationRangesÚIterationRangesEntryÚIterationRangesRootÚ
SIMDKernelÚSIMDScheduling)Ú	config_ofÚequal_1_arg_indicesÚnon_constexpr_signatureÚshould_unwrap_unspec_argÚsignature_to_meta)ÚSymbolicCallArg)Ú
ModuleType)ÚTypeVar©ÚDtypePropagationOpsHandler)ÚIRNode)ÚSIMDKernelFeaturesÚ_TÚ
perf_hintsÚscheduleÚfusionc                  ó@   — e Zd ZU dZi Zded<   i Zded<   edd„«       Zy)	ÚOpDtypeSupportzŸ
    Some Triton ops such as libdevice and tl.math only support float32 and float64.
    This class records which dtypes are supported by specific IR ops.
    z"dict[str, OrderedSet[torch.dtype]]Úsupported_dtypeszdict[str, bool]Úconvert_outputsc                ó¦   — |j                   }t        t        j                  t        j                  g«      | j
                  |<   || j                  |<   y ©N)Ú__name__r   ÚtorchÚfloat32Úfloat64rn   ro   )ÚclsÚfuncÚconvert_outputÚop_names       ú\/var/www/skyplay_api_hub/venv/lib/python3.12/site-packages/torch/_inductor/codegen/triton.pyÚregister_upcastzOpDtypeSupport.register_upcast   s=   € à—-‘-ˆÜ(2´E·M±MÄ5Ç=Á=Ð3QÓ(Rˆ×Ñ˜WÑ%Ø'5ˆ×Ñ˜GÒ$ó    N)rw   zCallable[..., str]rx   ÚboolÚreturnÚNone)	rr   Ú
__module__Ú__qualname__Ú__doc__rn   Ú__annotations__ro   Úclassmethodr{   © r|   rz   rm   rm   v   s1   … ñð
 <>ÐÐ8Ó=Ø')€O�_Ó)àò6ó ñ6r|   rm   c                 ód   — t        «       syddl} t        | j                  j                  d«      ryy)zd
    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    rz   Úgen_attr_descriptor_importr�   †   s-   € ô ÔØã#ô ˆv�‰×'Ñ'Ð):Ô;ØEàr|   c                 ó´   — t        «       } | j                  d«       t        «       x}r| j                  |«       | j                  d«       | j	                  «       S )NzD
        import triton
        import triton.language as tl
        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
        )rM   Úsplicer�   Ú	writelineÚgetvalue)ÚimportsÚ	attr_descs     rz   Úgen_common_triton_importsr”   ™   s[   € äÓ€GØ‡N�Nð	ôô /Ó0Ð0€yÐ0Ø×Ñ˜)Ô$à‡N�Nð	ôð ×ÑÓÐr|   c                  óÞ  — e Zd ZdZ eej                  ej                  g«      Z eej                  ej                  ej                  ge¢«      ZeD � ���ci c]%  }|t        j                  t        |   › d�dd¬«      “Œ' c}}}} ZeD � ���ci c]3  }|t        j                  t        |   j#                  «       › d�dd¬«      “Œ5 c}}}} Zed
d„«       Zed
d„«       Zy	c c}}}} w c c}}}} w )ÚTritonSymbolszU
    Stores sympy.Symbol instances and constants associated with triton codegen.
    ÚoffsetT©ÚintegerÚnonnegativeÚBLOCK©r™   Úpositivec                ó4   — | j                   |j                     S rq   )Úblock_sizesÚsymt©rv   Útrees     rz   Úget_block_sizezTritonSymbols.get_block_sizeÃ   s   € à�‰˜tŸy™yÑ)Ð)r|   c                ó4   — | j                   |j                     S rq   )Úblock_offsetsr    r¡   s     rz   Úget_block_offsetzTritonSymbols.get_block_offsetÇ   s   € à× Ñ  §¡Ñ+Ð+r|   N)r¢   rW   r~   zsympy.Symbol)rr   r€   r�   r‚   r   r   ÚR0_INDEXÚR1_INDEXÚreduction_typesÚXBLOCKÚYBLOCKÚZBLOCKÚblock_typesÚsympyÚSymbolr   r¥   ÚupperrŸ   r„   r£   r¦   )Ú.0r    r®   r   s   0000rz   r–   r–   ¯   s  „ ññ ! $§-¡-°·±Ð!?Ó@€OÙ˜dŸk™k¨4¯;©;¸¿¹ÐVÀoÐVÓW€Kð  ÷ñ àð 	Œe�l‰lœj¨Ñ.Ð/¨vÐ6ÀÐRVÔWÑWõ€Mð  ÷	ñ ð ð 	Œe�l‰lÜ˜$Ñ×%Ñ%Ó'Ð(¨Ð.¸Àtô
ñ 	
õ€Kð ò*ó ð*ð ò,ó ñ,ùõ#ùõ
s   Á*C
Â8C'
r–   c                  óz   — e Zd ZU ded<   ded<   ded<   ded<   d	ed
<   dd„Zdd„Zdd„Zdd„Zdd„Ze	dd„«       Z
y)ÚIndexingOptionsÚstrÚ	index_strúOrderedSet[str]Ú	mask_varsúOptional[str]Ú
expand_strr}   Ú_has_rindexú
sympy.ExprÚindexc                ó,   — t        | j                  «      S rq   )r}   r·   ©Úselfs    rz   Úhas_maskzIndexingOptions.has_maskÔ   s   € Ü�D—N‘NÓ#Ð#r|   c                óJ   — t        | j                  t        j                  «      S rq   )r   r¼   r   ÚTMPr¾   s    rz   Úhas_indirectzIndexingOptions.has_indirect×   s   € Ü" 4§:¡:¬t¯x©xÓ8Ð8r|   c                ó   — | j                   S rq   )rº   r¾   s    rz   Ú
has_rindexzIndexingOptions.has_rindexÚ   s   € Ø×ÑÐr|   c                ó:   — t        d„ | j                  D «       «      S )Nc              3  óP   K  — | ]  }t        |«      j                  d «      –— Œ  y­w)ÚtmpN©r´   Ú
startswith©r±   Úmasks     rz   ú	<genexpr>z.IndexingOptions.has_tmpmask.<locals>.<genexpr>Þ   s   è ø€ ÒJ°4”3�t“9×'Ñ'¨×.ÑJùó   ‚$&©Úanyr·   r¾   s    rz   Úhas_tmpmaskzIndexingOptions.has_tmpmaskÝ   s   € ÜÑJ¸4¿>¹>ÔJÓJÐJr|   c                ó:   — t        d„ | j                  D «       «      S )Nc              3  óP   K  — | ]  }t        |«      j                  d «      –— Œ  y­w)ÚrNrÉ   rË   s     rz   rÍ   z,IndexingOptions.has_rmask.<locals>.<genexpr>á   s   è ø€ ÒH°”3�t“9×'Ñ'¨×,ÑHùrÎ   rÏ   r¾   s    rz   Ú	has_rmaskzIndexingOptions.has_rmaskà   s   € ÜÑH¸¿¹ÔHÓHÐHr|   c                óp   — | j                   r)dj                  t        t        | j                   «      «      S dS )Nú & r   )r·   ÚjoinÚmapr´   r¾   s    rz   Úmask_strzIndexingOptions.mask_strã   s'   € à7;·~²~ˆu�z‰zœ#œc 4§>¡>Ó2Ó3ÐQÈ6ÐQr|   N©r~   r}   ©r~   r´   )rr   r€   r�   rƒ   rÀ   rÃ   rÅ   rÑ   rÕ   ÚpropertyrÚ   r…   r|   rz   r³   r³   Ì   sP   … àƒNØÓØÓØÓØÓó$ó9ó óKóIð òRó ñRr|   r³   c                  óR  — e Zd ZU ded<   ded<   ded<   ded<   d	ed
<   ded<   d	ed<   dZded<   ed!d„«       Zed!d„«       Zed!d„«       Zed!d„«       Z		 	 	 	 	 	 	 	 	 	 d"d„Z
e	 	 	 	 	 	 	 	 	 	 	 	 d#d„«       Z	 	 	 	 	 	 	 	 d$d„Zd%d&d„Zd'd„Zd(d„Zd)d„Zd*d„Zd*d„Zd*d„Zd*d„Zd*d „Zy)+ÚBlockPtrOptionsÚBlockParametersÚparamsr»   Úconstant_offsetú	list[int]Úorderr¶   r·   úSequence[sympy.Expr]Úbroadcast_shapez
list[bool]Úbroadcasting_dimsÚfinal_shapeNzOptional[list[int]]Ú_boundary_checkc                ó.   — | j                   j                  S rq   )rá   Úshaper¾   s    rz   rë   zBlockPtrOptions.shapeó   s   € à�{‰{× Ñ Ð r|   c                ó.   — | j                   j                  S rq   )rá   Úblock_shaper¾   s    rz   rí   zBlockPtrOptions.block_shape÷   s   € à�{‰{×&Ñ&Ð&r|   c                ó.   — | j                   j                  S rq   )rá   Ústridesr¾   s    rz   rï   zBlockPtrOptions.stridesû   ó   € à�{‰{×"Ñ"Ð"r|   c                ó.   — | j                   j                  S rq   )rá   Úoffsetsr¾   s    rz   rò   zBlockPtrOptions.offsetsÿ   rð   r|   c                ó0  ‡	— t        | j                  | j                  «      D ��cg c]#  \  }}|rt        j                  j
                  n|‘Œ% }}}t        |||«      }t        j                  j                  Š	t        | j                  «      xr7 t        |«      t        |«      k7  xs t        ˆ	fd„t        ||«      D «       «      }|r|r0d|› dt        j                  j                  | j                  «      › d�}t        || j                  |«      }|S c c}}w )z 
        Generate a broadcast and a reshape for the block pointer.
        This restores stride-0 dimensions which were removed from the block pointer.
        c              3  ór   •K  — | ].  \  }}‰j                  |d «      xs ‰j                  ||«       –— Œ0 y­w©rE   N)Ústatically_known_equals)r±   Úpre_dimÚpost_dimÚsizevarss      €rz   rÍ   z@BlockPtrOptions.codegen_broadcast_and_reshape.<locals>.<genexpr>  sG   øè ø€ ò ñ
 &�G˜Xð ×4Ñ4°W¸aÓ@ò KØ×7Ñ7¸ÀÓJôñùs   ƒ47útl.broadcast_to(ú, ú))Úzipræ   rç   r®   ÚSÚOneÚtriton_reshaperC   Úgraphrù   rÐ   ÚlenÚkernelÚindex_to_str)
r¿   ÚvalueÚinitial_shaperè   Úallow_implicitÚdimÚis_broadcastingÚpre_broadcast_shapeÚrequire_broadcastrù   s
            @rz   Úcodegen_broadcast_and_reshapez-BlockPtrOptions.codegen_broadcast_and_reshape  s  ø€ ô ),Ø×$Ñ$ d×&<Ñ&<ó)÷
á$��_ñ +ŒE�G‰G�KŠK°Ñ3ð
Ðñ 
ô ˜u mÐ5HÓIˆô —7‘7×#Ñ#ˆÜ × 6Ñ 6Ó7ò 	
ÜÐ#Ó$¬¨KÓ(8Ñ8ò Üó ô
 *-Ð-@À+Ó)Nôó ð 	ñ Ñ!2Ø& u g¨R´·±×0EÑ0EÀd×FZÑFZÓ0[Ð/\Ð\]Ð^ˆEô ˜u d×&:Ñ&:¸KÓHˆàˆùó=
s   ¥(Dc                óþ  ‡‡— t         j                  j                  Šd	ˆfd„} || j                  «      | _         || j                  «      | _        | j                  D �cg c]  }‰j                  |d«      ‘Œ }}| j                  D �cg c]  }‰j                  |d«      ‘Œ }	}t        |	«      rd|	d<   t        | j                  |	«      D ��
cg c]	  \  }}
|
s|‘Œ }}}
t        |	|«      D �cg c]  }t        |«      ‘Œ c}Šˆfd„}t        d
i t        j                  | «      j                  «       D ��ci c]  \  }}| ||«      “Œ c}}¤Ž} |D �cg c]  }t        j                  |«      ‘Œ }}t         j                   j"                  r%|d   j$                  dk(  sJ ‚|j'                  d«       t         j                   j(                  }t         j                   j*                  s…t-        | j                  «      t-        t         j                   j.                  «      |z
  k(  rIt         j                   j0                  j3                  «       r!|t4        j6                  j8                  g|z  z  }t;        | t         j                  j                  j=                  |«      t?        tA        tC        t-        | j                  «      «      «      «      ||||¬«      }|jE                  |«       |S c c}w c c}w c c}
}w c c}w c c}}w c c}w )z,Helper to create a  BlockPtrOptions instancec                óL   •— | D �cg c]  }‰j                  |«      ‘Œ c}S c c}w rq   )Úlookup_precomputed_size)ÚexprsÚexprrù   s     €rz   Úlookup_sizez+BlockPtrOptions.create.<locals>.lookup_size=  s"   ø€ ØGLÖM¸t�H×4Ñ4°TÕ:ÒMÐMùÒMs   †!r   rE   Féÿÿÿÿc                óR   •— t        | ‰«      D ��cg c]	  \  }}|s|‘Œ c}}S c c}}w )z@Removes any broadcasting or singleton dims from a given sequence)rý   )ÚitÚitemÚis_removableÚremovable_dimss      €rz   Úremove_dimsz+BlockPtrOptions.create.<locals>.remove_dims_  s3   ø€ ô +.¨b°.Ó*A÷á&�D˜,Ù#ò óð ùó s   ‘#Úx)rá   râ   rä   r·   rè   ræ   rç   )r  zIterable[sympy.Expr]r~   úlist[sympy.Expr]r…   )#rC   r  rù   rë   rï   rö   rí   Úallrý   rÐ   rà   ÚdataclassesÚasdictÚitemsr–   r£   r  Úno_x_dimÚprefixÚpopÚnum_reduction_dimsÚinside_reductionr  ÚnumelsÚfeaturesÚis_reductionr®   rþ   rÿ   rß   r  ÚlistÚreversedÚrangeÚcompute_boundary_check)rá   râ   Úrange_treesr·   Úget_max_blockr  Ústriderç   r  Úsingleton_dimsÚis_singletonræ   Údimsr  ÚkeyÚvalr¢   rè   Úreduction_ndimÚresultr  rù   s                       @@rz   ÚcreatezBlockPtrOptions.create0  sŽ  ù€ ô —7‘7×#Ñ#ˆõ	Nñ # 6§<¡<Ó0ˆŒÙ$ V§^¡^Ó4ˆŒð
 GMÇnÁnö
Ø<BˆH×,Ñ,¨V°QÕ7ð
Ðð 
ð AG×@RÑ@Rö
Ø9<ˆH×,Ñ,¨S°!Õ4ð
ˆð 
ô ˆ~Ôà!&ˆN˜2Ñô &)¨×);Ñ);¸^Ó%L÷
á!��\Ùò ð
ˆñ 
ô 14°NÐDUÓ0VÖW¨œ#˜d�)ÒWˆô	ô !ñ 
Ü5@×5GÑ5GÈÓ5O×5UÑ5UÓ5W×X©¨¨cˆs‘K Ó$Ñ$ÓXñ
ˆð
 GRÖR¸d”}×3Ñ3°DÕ9ÐRˆÐRÜ�8‰8×ÒØ˜q‘>×(Ñ(¨CÒ/Ð/Ð/Ø�O‰O˜AÔäŸ™×4Ñ4ˆä—‘×)Ò)Ü�F—N‘NÓ#¤s¬1¯8©8¯?©?Ó';¸nÑ'LÒLÜ—‘×!Ñ!×.Ñ.Ô0ð œEŸG™GŸK™K˜=¨>Ñ9Ñ9ˆKä ØÜŸG™G×,Ñ,×DÑDÀ_ÓUÜ”x¤¤c¨&¯,©,Ó&7Ó 8Ó9Ó:ØØ#Ø+Ø/ô
ˆð 	×%Ñ% mÔ4Øˆùò
ùò
ùó
ùò Xùó Yùò Ss$   ÁKÂK$ÃK)Ã+K/Ä2K4ÅK:c                óD   — t         j                  |   }t        |||i«      S )zN
        Replaces instances of {symt}_offset with the new expression.
        )r–   r¥   r<   )r¿   r  Úreplacementr    Úroffsets        rz   Úreplace_offsetzBlockPtrOptions.replace_offset‡  s&   € ô  ×-Ñ-¨dÑ3ˆÜ˜$ ¨+Ð 6Ó7Ð7r|   c           	     óÊ  ‡ — dˆ fd„}t         j                  j                  }g ‰ j                  ¢}|s|D �cg c]
  } ||«      ‘Œ }}‰ j                  dk7  r|› d |‰ j                  «      › d�n|d |‰ j
                  «      › �d |‰ j                  «      › �d |‰ j                  «      › �d |‰ j                  «      › �d	 ||«      › �g}d
dj                  |«      › d�S c c}w )a  
        Codegen a call to tl.make_block_ptr()

        Args:
            name: variable name for pointer
            roffset: should rn_offset be included in offsets=..., for use with tl.advance()

        Returns:
            "tl.make_block_ptr(...)"
        c                ó~   •— t         j                  D ](  }‰j                  | t        j                  d«      |«      } Œ* | S ©Nr   )r–   r©   r:  r®   ÚInteger)r  r    r¿   s     €rz   Úremove_roffsetsz/BlockPtrOptions.format.<locals>.remove_roffsetsœ  s<   ø€ Ü%×5Ñ5ò I�Ø×*Ñ*¨4´·±¸qÓ1AÀ4ÓH‘ðIàˆKr|   r   ú + (rü   zshape=zstrides=zblock_shape=zorder=zoffsets=ztl.make_block_ptr(rû   )r  r»   r~   r»   )
rC   r  r  rò   râ   rë   rï   rí   rä   rØ   )r¿   Únamer9  r?  Úfrò   r—   Úargss   `       rz   ÚformatzBlockPtrOptions.format�  sô   ø€ õ	ô
 �H‰H×!Ñ!ˆØ!�D—L‘L�/ˆÙØ=DÖE°6‘ vÕ.ÐEˆGÐEð ×'Ñ'¨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ùò Fs   ¶C c           
     óô  — t         j                  j                  }t        j                  j                  «       D ��ci c]  \  }}| |t        |   «      “Œ }}}t        t        | j                  «      «      D �cg c]è  }|j                  | j                  |   t        j                  j                  «      s¯|j                  | j                  |   | j                   |   «      sƒ|j                  | j                  |   t#        | j                   |   |«      «      sMt         j$                  j&                  r1| j                   |   t        j                  t(        j*                     k(  s|‘Œê c}| _        yc c}}w c c}w )z6List of indices to pass to tl.load(boundary_check=...)N)rC   r  rù   r–   rŸ   r  r   r*  r  rë   rö   rï   r®   rþ   ÚZeroÚstatically_known_multiple_ofrí   r<   r  r   r   rª   ré   )r¿   r-  rù   r    Ú
block_sizeÚblock_to_maxÚidxs          rz   r+  z&BlockPtrOptions.compute_boundary_check³  s-  € ä—7‘7×#Ñ#ˆô %2×$=Ñ$=×$CÑ$CÓ$E÷/
á ��jð ™¤j°Ñ&6Ó7Ñ7ð/
ˆñ /
ô œS §¡›_Ó-ö 
àà×4Ñ4°T·\±\À#Ñ5FÌÏÉÏÉÔUØ ×=Ñ=Ø—J‘J˜s‘O T×%5Ñ%5°cÑ%:ôð !×=Ñ=Ø—J‘J˜s‘O¤Z°×0@Ñ0@ÀÑ0EÀ|Ó%Tôô —H‘H×%Ò%Ø×(Ñ(¨Ñ-´×1JÑ1JÌ4Ï;É;Ñ1WÒWò ò 
ˆÕùó/
ùò
 
s   ¼E/Á9C-E5c                ó6   — | j                   €J ‚| j                   S rq   )ré   r¾   s    rz   Úboundary_checkzBlockPtrOptions.boundary_checkÐ  s    € Ø×#Ñ#Ð/Ð/Ð/Ø×#Ñ#Ð#r|   c           	     óâ   — t         j                  |   }| j                  D �cg c]A  }| j                  |||«      | j                  |t        j
                  j                  |«      z
  ‘ŒC }}|S c c}w )av  
        Codegen string to pass to tl.advance(name, ...).

        Advance is the difference between offsets in each loop iteration.
        To compute it, we replace rN_offset with multiples of RN_BLOCK.
        Since we expect rN_offset to vary in range(0, rN_numel, RN_BLOCK), the first
        iteration has rN_offset=0, while the second has rN_offset=RN_BLOCK.
        )r–   rŸ   rò   r:  r®   rþ   rF  )r¿   r    Úrblockr—   Úadvances        rz   Úadvance_roffsetzBlockPtrOptions.advance_roffsetÔ  st   € ô ×*Ñ*¨4Ñ0ˆð Ÿ,™,ö
ð
 ð ×#Ñ# F¨F°DÓ9Ø×%Ñ% f¬e¯g©g¯l©l¸DÓAóBð
ˆð 
ð ˆùò
s   ¢AA,c                 ó   — y©NFr…   r¾   s    rz   rÃ   zBlockPtrOptions.has_indirectç  ó   € Ør|   c                ó:   — t        d„ | j                  D «       «      S )Nc              3  óP   K  — | ]  }t        |t        j                  «      –— Œ  y ­wrq   )r   r–   r©   )r±   r  s     rz   rÍ   z-BlockPtrOptions.has_rindex.<locals>.<genexpr>ë  s%   è ø€ ò 
àô   ¤m×&CÑ&C×Dñ
ùrÎ   )rÐ   rí   r¾   s    rz   rÅ   zBlockPtrOptions.has_rindexê  s"   € Üñ 
à×(Ñ(ô
ó 
ð 	
r|   c                ó"   — | j                  «       S rq   )rÅ   r¾   s    rz   rÕ   zBlockPtrOptions.has_rmaskð  s   € Ø�‰Ó Ð r|   c                 ó   — yrR  r…   r¾   s    rz   rÑ   zBlockPtrOptions.has_tmpmaskó  rS  r|   c                ó4   — t        | j                  «       «      S rq   )r}   rL  r¾   s    rz   rÀ   zBlockPtrOptions.has_maskö  s   € Ü�D×'Ñ'Ó)Ó*Ð*r|   ©r~   r  )
r  r´   r  rå   rè   rå   r  r}   r~   r´   )rá   rà   râ   r»   r,  zlist[IterationRangesRoot]r·   r¶   r-  úCallable[[str], int]r~   rß   )r  r»   r8  r»   r    r   r~   r»   ©T)rA  r´   r~   r´   )r-  rZ  r~   r   )r~   rã   )r    r   r~   r»   rÛ   )rr   r€   r�   rƒ   ré   rÝ   rë   rí   rï   rò   r  Ústaticmethodr6  r:  rD  r+  rL  rP  rÃ   rÅ   rÕ   rÑ   rÀ   r…   r|   rz   rß   rß   è   se  … àÓØÓØÓØÓØ)Ó)Ø!Ó!Ø%Ó%Ø+/€OÐ(Ó/àò!ó ð!ð ò'ó ð'ð ò#ó ð#ð ò#ó ð#ð+àð+ð ,ð+ð *ð	+ð
 ð+ð 
ó+ðZ ðTàðTð $ðTð /ð	Tð
 #ðTð ,ðTð 
òTó ðTðl8Øð8Ø-7ð8Ø?Cð8à	ó8ô!7óF
ó:$óó&ó
ó!óô+r|   rß   c                ór  — t        |t        «      rt        |t        «      sJ ‚|D �cg c]!  }t        j                  j	                  |«      ‘Œ# }}|D �cg c]!  }t        j                  j	                  |«      ‘Œ# }}||k(  r| S |D �cg c]
  }|dk7  sŒ	|‘Œ c}|k7  rd| › ddj                  |«      › d�S d}g }|D ]G  }	|t        |«      k  r|	||   k(  r|j                  d«       |dz  }Œ0|	dk(  sJ ‚|j                  d	«       ŒI |t        |«      k(  sJ ‚| › d
dj                  |«      › d�S c c}w c c}w c c}w )z7Workaround https://github.com/openai/triton/issues/2836Ú1ztl.reshape(z, [rû   z])r   ú:rE   r   ú[ú])Ú
isinstancer(  rC   r  r  rØ   r  Úappend)
r  Ú	old_shapeÚ	new_shaperë   Úold_shape_strÚnew_shape_strÚsrJ  ÚexpandÚsizes
             rz   r   r   ú  sE  € ô �i¤Ô&¬:°iÄÔ+FÐFÐFà?HÖI°e”Q—X‘X×*Ñ*¨5Õ1ÐI€MÐIØ?HÖI°e”Q—X‘X×*Ñ*¨5Õ1ÐI€MÐIà˜Ò%ØˆØ Ö-ˆa 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   §&D*Á&D/Â
D4ÂD4c                  óð   — e Zd Zdd„Zdd„Zdd„Zdd„Zdd„Zdd„Zdd„Z	dd„Z
dd	„Zdd
„Zdd„Zdd„ZeZdd„Zdd„Zdd„Zdd„Zdd„Zdd„Zdd„Zdd„Zdd„Zdd„Zdd„Zdd„Zdd„Zdd„Zdd„Zdd„Zy) ÚTritonPrinterc                ó°   — t        |j                  «      dk(  sJ ‚d| j                  |j                  d   «      › dt        j                  j
                  › d�S )NrE   úlibdevice.trunc(r   ú).to(rü   ©r  rC  Ú_printrC   r  Úindex_dtype©r¿   r  s     rz   Ú_print_TruncToIntzTritonPrinter._print_TruncToInt  óM   € Ü�4—9‘9‹~ Ò"Ð"Ð"à˜tŸ{™{¨4¯9©9°Q©<Ó8Ð9¸¼q¿x¹x×?SÑ?SÐ>TÐTUÐVð	
r|   c                óx   — t        j                  «       rt        j                  j                  r|› }|S d|› d�}|S )Nztl.full([], z, tl.float64))r   Ú	is_fbcoders   ÚversionÚhip)r¿   r  Úrets      rz   Ú_print_FloatzTritonPrinter._print_Float   s=   € Ü×ÑÔ¤%§-¡-×"3Ò"3Ø�FˆCð ˆ
ð !   mÐ4ˆCØˆ
r|   c                ó’   — t        |j                  «      dk(  sJ ‚| j                  |j                  d   t        d   dz
  «      }|› d�S )NrE   r   ÚAtomç      à?z.to(tl.float64))r  rC  Úparenthesizer   )r¿   r  rh  s      rz   Ú_print_ToFloatzTritonPrinter._print_ToFloat'  sI   € Ü�4—9‘9‹~ Ò"Ð"Ð"Ø×Ñ˜dŸi™i¨™l¬J°vÑ,>ÀÑ,DÓEˆØ��OÐ$Ð$r|   c                óô   — |j                   \  }}|j                  r3|j                  r'| j                  |j                   dt        d   dz
  «      S | j	                  |«      }| j	                  |«      }d|› d|› d�S )Nú % r}  r~  z!triton_helpers.remainder_integer(rû   rü   )rC  Úis_nonnegativeÚ	stringifyr   rq  ©r¿   r  ÚquotÚdivÚquot_sÚdiv_ss         rz   Ú_print_PythonModzTritonPrinter._print_PythonMod,  sr   € Ø—I‘I‰	ˆˆcØ×Ò 3×#5Ò#5Ø—>‘> $§)¡)¨U´J¸vÑ4FÈÑ4LÓMÐMØ—‘˜TÓ"ˆØ—‘˜CÓ ˆØ2°6°(¸"¸U¸GÀ1ÐEÐEr|   c                ó  — |j                   sJ ‚|j                  \  }}|j                  r3|j                  r'| j                  |j                  dt        d   dz
  «      S | j                  |«      }| j                  |«      }d|› d|› d�S )Nú // r}  r~  z!triton_helpers.div_floor_integer(z,  rü   )Ú
is_integerrC  rƒ  r„  r   rq  r…  s         rz   Ú_print_FloorDivzTritonPrinter._print_FloorDiv4  s~   € Ø�ŠÐˆØ—I‘I‰	ˆˆcØ×Ò 3×#5Ò#5Ø—>‘> $§)¡)¨V´ZÀÑ5GÈ#Ñ5MÓNÐNØ—‘˜TÓ"ˆØ—‘˜CÓ ˆØ2°6°(¸#¸e¸WÀAÐFÐFr|   c                óP   — | j                  |j                  dt        d   dz
  «      S )Nú / r}  r~  )r„  rC  r   rs  s     rz   Ú_print_IntTrueDivzTritonPrinter._print_IntTrueDiv?  s#   € Ø�~‰~˜dŸi™i¨´
¸6Ñ0BÀSÑ0HÓIÐIr|   c                ó°   — t        |j                  «      dk(  sJ ‚d| j                  |j                  d   «      › dt        j                  j
                  › d�S ©NrE   úlibdevice.floor(r   ro  rü   rp  rs  s     rz   Ú_print_floorzTritonPrinter._print_floorD  ru  r|   c                ó°   — t        |j                  «      dk(  sJ ‚d| j                  |j                  d   «      › dt        j                  j
                  › d�S r“  rp  rs  s     rz   Ú_print_FloorToIntzTritonPrinter._print_FloorToIntJ  ru  r|   c                ó°   — t        |j                  «      dk(  sJ ‚d| j                  |j                  d   «      › dt        j                  j
                  › d�S ©NrE   úlibdevice.ceil(r   ro  rü   rp  rs  s     rz   Ú_print_ceilingzTritonPrinter._print_ceilingP  óK   € Ü�4—9‘9‹~ Ò"Ð"Ð"Ø  §¡¨T¯Y©Y°q©\Ó!:Ð ;¸5ÄÇÁ×AUÑAUÐ@VÐVWÐXÐXr|   c                ó°   — t        |j                  «      dk(  sJ ‚d| j                  |j                  d   «      › dt        j                  j
                  › d�S r™  rp  rs  s     rz   Ú_print_CeilToIntzTritonPrinter._print_CeilToIntT  rœ  r|   c                ó,   — d| j                  |«      › d�S )Nzlibdevice.sqrt((ú).to(tl.float32)))rq  rs  s     rz   Ú_helper_sqrtzTritonPrinter._helper_sqrtX  s   € Ø! $§+¡+¨dÓ"3Ð!4Ð4EÐFÐFr|   c                ó„   — d| j                  |j                  d   «      › d| j                  |j                  d   «      › d�S )Núlibdevice.pow(r   rû   rE   rü   )rq  rC  rs  s     rz   Ú_print_FloatPowzTritonPrinter._print_FloatPow[  s?   € à˜TŸ[™[¨¯©°1©Ó6Ð7°r¸$¿+¹+ÀdÇiÁiÐPQÁlÓ:SÐ9TÐTUÐVð	
r|   c                óÎ   — | j                  |j                  d   «      }| j                  |j                  d   «      }| j                  |j                  d   «      }d|› d|› d|› d�S )Nr   rE   r   ú	tl.where(rû   rü   )ÚdoprintrC  )r¿   r  ÚcÚpÚqs        rz   Ú_print_WherezTritonPrinter._print_Whereb  s_   € Ø�L‰L˜Ÿ™ 1™Ó&ˆØ�L‰L˜Ÿ™ 1™Ó&ˆØ�L‰L˜Ÿ™ 1™Ó&ˆØ˜1˜#˜R ˜s " Q C qÐ)Ð)r|   c                óº  — t        |j                  «      dk(  r| j                  |j                  d   «      S t        |j                  «      dz  }t        |«      }| j                   ||j                  d| Ž «      }| j                   ||j                  |d Ž «      }t	        d„ ||fD «       «      \  }}|dv sJ d|› d�«       ‚d	|› d
|› d|› d|› d|› d
|› d|› d|› d�S )zK
        Helper for max/min code genereration.
        cmp: > or <
        rE   r   r   Nc              3  ó(   K  — | ]
  }d |› d�–— Œ y­w)ú(rü   Nr…   ©r±   r  s     rz   rÍ   z6TritonPrinter._print_min_max_helper.<locals>.<genexpr>w  s   è ø€ Ò. !�q˜˜˜1”XÑ.ùs   ‚)ú>ú<zUnexpected comparator: 'ú'r®  z * (ú z= z) + ú)))r  rC  rq  ÚtypeÚtuple)r¿   r  ÚcmpÚmidrv   ÚaÚbs          rz   Ú_print_min_max_helperz#TritonPrinter._print_min_max_helperh  sð   € ô
 ˆt�y‰y‹>˜QÒØ—;‘;˜tŸy™y¨™|Ó,Ð,ä�$—)‘)‹n Ñ!ˆÜ�4‹jˆØ�K‰K™˜TŸY™Y t¨˜_Ð-Ó.ˆØ�K‰K™˜TŸY™Y s t˜_Ð-Ó.ˆô Ñ.¨¨1 vÔ.Ó.‰ˆˆ1Ø�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                ó&   — | j                  |d«      S )Nr±  ©r»  rs  s     rz   Ú
_print_MinzTritonPrinter._print_Min{  ó   € Ø×)Ñ)¨$°Ó4Ð4r|   c                ó&   — | j                  |d«      S )Nr°  r½  rs  s     rz   Ú
_print_MaxzTritonPrinter._print_Max~  r¿  r|   c                óz   — t        |j                  «      dk(  sJ ‚d| j                  |j                  d   «      › d�S )NrE   útl_math.abs(r   rü   ©r  rC  rq  rs  s     rz   Ú
_print_AbszTritonPrinter._print_Abs�  s9   € Ü�4—9‘9‹~ Ò"Ð"Ð"Ø˜dŸk™k¨$¯)©)°A©,Ó7Ð8¸Ð:Ð:r|   c                óz   — t        |j                  «      dk(  sJ ‚d| j                  |j                  d   «      › d�S )NrE   zlibdevice.cos((r   r   rÄ  rs  s     rz   Ú_print_OpaqueUnaryFn_cosz&TritonPrinter._print_OpaqueUnaryFn_cos…  ó:   € Ü�4—9‘9‹~ Ò"Ð"Ð"Ø  §¡¨T¯Y©Y°q©\Ó!:Ð ;Ð;LÐMÐMr|   c                óz   — t        |j                  «      dk(  sJ ‚d| j                  |j                  d   «      › d�S )NrE   zlibdevice.cosh((r   r   rÄ  rs  s     rz   Ú_print_OpaqueUnaryFn_coshz'TritonPrinter._print_OpaqueUnaryFn_cosh‰  ó:   € Ü�4—9‘9‹~ Ò"Ð"Ð"Ø! $§+¡+¨d¯i©i¸©lÓ";Ð!<Ð<MÐNÐNr|   c                óz   — t        |j                  «      dk(  sJ ‚d| j                  |j                  d   «      › d�S )NrE   zlibdevice.acos((r   r   rÄ  rs  s     rz   Ú_print_OpaqueUnaryFn_acosz'TritonPrinter._print_OpaqueUnaryFn_acos�  rË  r|   c                óz   — t        |j                  «      dk(  sJ ‚d| j                  |j                  d   «      › d�S )NrE   zlibdevice.sin((r   r   rÄ  rs  s     rz   Ú_print_OpaqueUnaryFn_sinz&TritonPrinter._print_OpaqueUnaryFn_sin‘  rÈ  r|   c                óz   — t        |j                  «      dk(  sJ ‚d| j                  |j                  d   «      › d�S )NrE   zlibdevice.sinh((r   r   rÄ  rs  s     rz   Ú_print_OpaqueUnaryFn_sinhz'TritonPrinter._print_OpaqueUnaryFn_sinh•  rË  r|   c                óz   — t        |j                  «      dk(  sJ ‚d| j                  |j                  d   «      › d�S )NrE   zlibdevice.asin((r   r   rÄ  rs  s     rz   Ú_print_OpaqueUnaryFn_asinz'TritonPrinter._print_OpaqueUnaryFn_asin™  rË  r|   c                óz   — t        |j                  «      dk(  sJ ‚d| j                  |j                  d   «      › d�S )NrE   zlibdevice.tan((r   r   rÄ  rs  s     rz   Ú_print_OpaqueUnaryFn_tanz&TritonPrinter._print_OpaqueUnaryFn_tan�  rÈ  r|   c                óz   — t        |j                  «      dk(  sJ ‚d| j                  |j                  d   «      › d�S )NrE   zlibdevice.tanh((r   r   rÄ  rs  s     rz   Ú_print_OpaqueUnaryFn_tanhz'TritonPrinter._print_OpaqueUnaryFn_tanh¡  rË  r|   c                óz   — t        |j                  «      dk(  sJ ‚d| j                  |j                  d   «      › d�S )NrE   zlibdevice.atan((r   r   rÄ  rs  s     rz   Ú_print_OpaqueUnaryFn_atanz'TritonPrinter._print_OpaqueUnaryFn_atan¥  rË  r|   c                ó°   — t        |j                  «      dk(  sJ ‚d| j                  |j                  d   «      › dt        j                  j
                  › d�S )NrE   zlibdevice.llrint(r   ro  rü   rp  rs  s     rz   Ú_print_RoundToIntzTritonPrinter._print_RoundToInt©  sM   € Ü�4—9‘9‹~ Ò"Ð"Ð"à §¡¨D¯I©I°a©LÓ 9Ð:¸%ÄÇÁ×@TÑ@TÐ?UÐUVÐWð	
r|   c                óâ   — t        |j                  «      dk(  sJ ‚|j                  \  }}|j                  r|dk  sJ ‚t        d|› d�«      ‚| j	                  |t
        d   «      }d|› d|› d| › �S )	Nr   r   zOFor integer inputs, only non-negative ndigits are currently supported, but got ú.ÚMulzlibdevice.nearbyint(1eú * z) * 1e)r  rC  r�  Ú
ValueErrorr  r   )r¿   r  ÚnumberÚndigitsÚ
number_strs        rz   Ú_print_RoundDecimalz!TritonPrinter._print_RoundDecimal¯  sŠ   € Ü�4—9‘9‹~ Ò"Ð"Ð"ØŸ)™)‰ˆ�Ø×Òà˜Q’;Ð�;ÜØaÐbiÐajÐjkÐlóð ð ×&Ñ& v¬z¸%Ñ/@ÓAˆ
Ø'¨ y°°J°<¸vÀwÀhÀZÐPÐPr|   N)r  r»   r~   r´   )r  r»   r·  r´   r~   r´   ) rr   r€   r�   rt  r{  r€  rŠ  rŽ  r‘  r•  r—  r›  rž  r¡  r¤  Ú_print_PowByNaturalr«  r»  r¾  rÁ  rÅ  rÇ  rÊ  rÍ  rÏ  rÑ  rÓ  rÕ  r×  rÙ  rÛ  rä  r…   r|   rz   rl  rl    s§   „ ó
óó%ó
FóGóJó

ó
óYóYóGó
ð
 *Ðó*óCó&5ó5ó;óNóOóOóNóOóOóNóOóOó
ôQr|   rl  c                ó*   — t        t        | «      «      S )zCConvert torch.dtype to triton type and upcast [b]float16 to float32)r=   r?   ©Údtypes    rz   Útriton_compute_typeré  À  s   € äÔ*¨5Ó1Ó2Ð2r|   c                óB   — t        | «      } t        | dd«      }|r|dz  S y)z'Number of bits of triton_compute_type()ÚitemsizeNé   r  )r?   Úgetattr)rè  rë  s     rz   Ú_get_primitive_bitwidthrî  Å  s*   € ä Ó&€EÜ�u˜j¨$Ó/€HÙØ˜!‰|Ðàr|   c                ó^   — | t         j                  k(  rt         j                  } t        | «      S )z@Convert torch.dtype to triton type, with fix for storing tl.bool)rs   r}   Úint8r=   rç  s    rz   Útriton_store_typerñ  Ï  s"   € à”—
‘
ÒÜ—
‘
ˆÜ�uÓÐr|   c                ó„   — t        | «      r+| j                  r| j                  dk  rt        j                  S t        | «      S )z0Implicit upcasts used for Triton reduction typesé   )r   Ú	is_signedrë  rs   Úint32r?   rç  s    rz   Úupcast_acc_dtyperö  Ö  s0   € ä˜Ô 5§?¢?°u·~±~ÈÒ7JÜ�{‰{ÐÜ˜uÓ%Ð%r|   c                ó*   — t        t        | «      «      S )z:Convert torch.dtype to triton type, with reduction upcasts)ré  rö  rç  s    rz   Útriton_acc_typerø  Ý  s   € äÔ/°Ó6Ó7Ð7r|   c                ó<   — | j                   dk  xr | j                  S )Nr   )rë  Úis_floating_pointrç  s    rz   Úlow_precision_fprû  â  s   € Ø�>‰>˜QÑÒ: 5×#:Ñ#:Ð:r|   c                óŠ   — t        | t        «      sy| j                  }t        |t        j                  «      rt	        |«      S dS rR  )rb  rK   rè  rs   rû  )Úvarrè  s     rz   Úlow_precision_fp_varrþ  æ  s6   € Ü�cœ;Ô'Øà�I‰I€EÜ&0°¼¿¹Ô&DÔ˜EÓ"ÐOÈ%ÐOr|   c                  ó&   ‡ — e Zd Zdˆ fd„Zd„ Zˆ xZS )ÚTritonCSEVariablec                óf   •— t         ‰| �  |||«       t        t           «       | _        |€J d«       ‚y )Nz!TritonCSEVariable must have dtype)ÚsuperÚ__init__r   r´   r·   )r¿   rA  Úboundsrè  Ú	__class__s       €rz   r  zTritonCSEVariable.__init__ï  s4   ø€ Ü‰Ñ˜˜v uÔ-ä#¤C™Ó*ˆŒØÐ ÐEÐ"EÓEÑ r|   c                óF  — |D ]œ  }t        |t        «      r&| j                  j                  |j                  «       Œ9t        |t        j
                  «      sŒTt        j                  D ]6  }t        ||«      sŒ| j                  j                  t        |   › d�g«        Œœ Œž y )NrÌ   )
rb  r   r·   Úupdater®   r¯   r–   r­   r   r   )r¿   rA  rC  ÚkwargsÚargr    s         rz   Úupdate_on_argsz TritonCSEVariable.update_on_argsõ  s†   € Øò 
	ˆCÜ˜#Ô0Ô1Ø—‘×%Ñ% c§m¡mÕ4Ü˜C¤§¡Õ.ô *×5Ñ5ò �DÜ% c¨4Õ0ØŸ™×-Ñ-´*¸TÑ2BÐ1CÀ4Ð/HÐ.IÔJÙññ
	r|   )r  zValueRanges[Any]rè  útorch.dtyper~   r   )rr   r€   r�   r  r
  Ú__classcell__©r  s   @rz   r   r   î  s   ø„ õFör|   r   c                 ó   — ddl m}   | «       S )Nr   rd   )Ú!torch._inductor.dtype_propagationre   rd   s    rz   Úget_dtype_handlerr    s   € ÝLá%Ó'Ð'r|   c                ó0   ‡ ‡‡— dd„Šdˆfd„Šdˆ ˆˆfd„}|S )z“
    Codegen helper to upcast arguments to float32, depending on the config and dtype.
    This decorates tl.math/libdevice codegen functions.
    c                ó¸   — t         j                  j                   xr> t        | t        «      xr, | j
                  t        j                  t        j                  fv S rq   )	r   rŒ   Úcodegen_upcast_to_fp32rb  rK   rè  rs   Úfloat16Úbfloat16)rý  s    rz   Úneeds_upcastz*maybe_upcast_float32.<locals>.needs_upcast  sD   € ä—‘×4Ñ4Ð4ò =Ü˜3¤Ó,ò=à—	‘	œeŸm™m¬U¯^©^Ð<Ð<ð	
r|   c                ó(   •—  ‰| «      rdnd}| › |› �S )Nú.to(tl.float32)r‡   r…   )rý  Úupcast_stringr  s     €rz   Úmaybe_upcast_argz.maybe_upcast_float32.<locals>.maybe_upcast_arg  s!   ø€ Ù-9¸#Ô->Ñ)ÀBˆØ��}�oÐ&Ð&r|   c                óH   •‡ — t         j                  ‰ ‰«       dˆˆ ˆˆfd„}|S )Nc                 óÈ  •— | D �cg c]
  } ‰|«      ‘Œ }}|j                  «       D ��ci c]  \  }}| ‰|«      “Œ }}} ‰|i |¤Ž}‰xr6 t        ˆfd„t        j                  | |j	                  «       «      D «       «      }|sd n# t        t        «       ‰j                  «      | i |¤Ž}	|	t        j                  d fv}
|
r|	�dt        |	«      › d�nd}|› |› �S c c}w c c}}w )Nc              3  ó.   •K  — | ]  } ‰|«      –— Œ y ­wrq   r…   )r±   rý  r  s     €rz   rÍ   zKmaybe_upcast_float32.<locals>.decorator.<locals>.wrapped.<locals>.<genexpr>%  s   øè ø€ ò 6Ø&)‘˜S×!ñ6ùs   ƒú.to(rü   r‡   )r  rÐ   Ú	itertoolsÚchainÚvaluesrí  r  rr   rs   rt   r=   )rC  r  r	  Úupcast_argsr2  r3  Úupcast_kwargsr5  Úany_needs_upcastÚresult_dtypeÚneeds_downcastÚdowncast_stringrx   rw   r  r  s               €€€€rz   Úwrappedz8maybe_upcast_float32.<locals>.decorator.<locals>.wrapped  s  ø€ à<@ÖA°SÑ+¨CÕ0ÐAˆKÐAØHNÏÉË×W¹H¸CÀ˜SÑ"2°3Ó"7Ñ7ÐWˆMÑWñ ˜;Ð8¨-Ñ8ˆFØ-ò  ´#ó 6Ü-6¯_©_¸TÀ6Ç=Á=Ã?Ó-Sô6ó 3Ðñ
 (ñ à@”WÔ.Ó0°$·-±-Ó@À$ÐQÈ&ÑQð ð
 *´%·-±-ÀÐ1FÐFˆNñ " lÐ&>ð ”{ <Ó0Ð1°Ñ3àð ð
 �X˜oÐ.Ð/Ð/ùò' BùÛWs
   †CªCrÜ   )rm   r{   )rw   r(  rx   r  r  s   ` €€€rz   Ú	decoratorz'maybe_upcast_float32.<locals>.decorator  s$   ù€ ä×&Ñ& t¨^Ô<÷	0ð 	0ð. ˆr|   rÛ   rÜ   )rw   úCallable[..., Any]r~   r*  r…   )rx   r)  r  r  s   ` @@rz   Úmaybe_upcast_float32r+  	  s   ú€ ó
õ'÷ð: Ðr|   c                  óÄ  — e Zd ZdZ ej
                  ej                  «      Ze	 	 dT	 	 	 dUd„«       Z	edVd„«       Z
ed„ «       Zed„ «       Ze e«       d„ «       «       Zed	„ «       Zed
„ «       Ze e«       d„ «       «       Ze e«       d„ «       «       Ze e«       d„ «       «       Ze e«       d„ «       «       Ze e«       d„ «       «       Ze e«       d„ «       «       Ze e«       d„ «       «       Zed„ «       Zed„ «       Zed„ «       Zed„ «       Zedej<                  dddœd„«       Ze e«       d„ «       «       Z e e«       d„ «       «       Z!e e«       d„ «       «       Z"e e«       d„ «       «       Z#ed„ «       Z$ed„ «       Z%e e«       d„ «       «       Z&e e«       d „ «       «       Z'e e«       d!„ «       «       Z(e e«       d"„ «       «       Z)e e«       d#„ «       «       Z*e e«       d$„ «       «       Z+e e«       d%„ «       «       Z,e e«       d&„ «       «       Z-e e«       d'„ «       «       Z.e e«       d(„ «       «       Z/e e«       d)„ «       «       Z0e e«       d*„ «       «       Z1e e«       d+„ «       «       Z2e e«       d,„ «       «       Z3e e«       d-„ «       «       Z4e e«       d.„ «       «       Z5e e«       d/„ «       «       Ze e«       d0„ «       «       Z6ed1„ «       Z7ed2„ «       Z8ed3„ «       Z9ed4„ «       Z:ed5„ «       Z;ed6„ «       Z<ed7„ «       Z=ed8„ «       Z>ed9„ «       Z?ed:„ «       Z@ed;„ «       ZAed<„ «       ZBed=„ «       ZCed>„ «       ZDe e«       d?„ «       «       ZEe e«       d@„ «       «       ZFe e«       dA„ «       «       ZGe e«       dB„ «       «       ZHe e«       dC„ «       «       ZIedD„ «       ZJe e«       dE„ «       «       ZKe e«       dF„ «       «       ZLe e«       dG„ «       «       ZMe e«       dH„ «       «       ZNe edI¬J«      dK„ «       «       ZOe edI¬J«      dL„ «       «       ZPe e«       dM„ «       «       ZQe e«       dN„ «       «       ZRedO„ «       ZSedP„ «       ZTe e«       dQ„ «       «       ZUedR„ «       ZVe e«       dS„ «       «       ZWy)WÚTritonOverrideszMap element-wise ops to TritonNTc                ó:  — 	 	 	 	 	 	 dd„}|�>t         |||«      t        j                  j                  «      t        j                  _        |t        j
                  k(  rd| › d�S |t        j                  k(  r| › d�S |rt        |«      }nt        |«      }| › d|› d�S )Nc                ó  — | |k(  ryt         j                  t         j                  f}| |v r||v r| |k7  rJ d«       ‚| t         j                  k(  s|t         j                  k(  ry| t         j                  k(  s|t         j                  k(  ryy)Nr   zCConversions between float8_e5m2 and float8_e4m3fn is not supported!ró  r   )rs   Úfloat8_e4m3fnÚfloat8_e5m2)Ú	src_dtypeÚ	dst_dtypeÚ
fp8_dtypess      rz   Ú_get_min_elements_per_threadz>TritonOverrides.to_dtype.<locals>._get_min_elements_per_threadF  sœ   € ð ˜IÒ%àô ×#Ñ#Ü×!Ñ!ðˆJð ˜ZÑ'Ø Ñ+Ø Ò*ðUð Uó	Uð ð
 œE×-Ñ-Ò-°¼e×>OÑ>OÒ1OØØœE×/Ñ/Ò/°9Ä×@SÑ@SÒ3SØàr|   r®  z != 0)z.to(tl.int8).to(tl.uint8)r  rü   )r2  r  r3  r  r~   Úint)	ÚmaxrC   r  Úmin_elem_per_threadrs   r}   Úuint8ré  rñ  )r  rè  r2  Úuse_compute_typesr5  Ú	out_dtypes         rz   Úto_dtypezTritonOverrides.to_dtype?  s³   € ð	Ø"ð	Ø/:ð	àó	ð6 Ð ô ,/Ù,¨Y¸Ó>Ü—‘×,Ñ,ó,ŒA�H‰HÔ(ð
 ”E—J‘JÒØ�q�c˜�=Ð Ø”e—k‘kÒ!ð �SÐ1Ð2Ð2áÜ+¨EÓ2‰Iä)¨%Ó0ˆIà��D˜˜ 1Ð%Ð%r|   c                ó  — t        |«      }|t        j                  t        j                  fv r´t        j
                  j                  ršt        |«      j                  d«      d   }| › d|› d�}|t        j                  t        j                  fv r"t        |«      j                  d«      d   }d|› �}|› d|› d�}|t        j                  t        j                  fv r|› d�S |S t        |«      }t        |«      }||k(  rd	nd
}	| › d|› d|	› d�S )NrÝ  r  z.to(tl.rü   ztl.r  z, bitcast=True)r  ÚTrueÚFalsez
, bitcast=)
ré  rs   r  r  r   rŒ   r  r´   Úsplitrî  )
r  rè  r2  Útriton_dtypeÚtriton_src_dtypeÚcast_xÚtriton_type_nameÚsrc_dtype_bitwidthÚtarget_dtype_bitwidthÚbitcasts
             rz   Úto_dtype_bitcastz TritonOverrides.to_dtype_bitcastx  s  € ä*¨5Ó1ˆð œ%Ÿ-™-¬¯©Ð8Ñ8Ü—‘×4Ò4ä" 9›~×3Ñ3°CÓ8¸Ñ<ÐØ�s˜'Ð"2Ð!3°1Ð5ˆFØœŸ™¬¯©Ð7Ñ7Ü#& u£:×#3Ñ#3°CÓ#8¸Ñ#<Ð Ø!$Ð%5Ð$6Ð7�Ø�x˜t L >°ÐAˆFØœŸ™¬¯©Ð7Ñ7Ø ˜ Ð1Ð1ØˆMä!8¸Ó!CÐÜ$;¸EÓ$BÐ!Ø 2Ð6KÒ K‘fÐQXˆGØ�S˜˜\˜N¨*°W°I¸QÐ?Ð?r|   c                óž   — t         j                  j                  |«      }t         || «      «      }t	        |«      }|dk(  r|S d|› d|› d|› d�S )Nz
tl.float32útl.full(rû   rü   )rs   Ú_prims_commonÚdtype_to_typerV   ré  )r  rè  rë   Útype_Ú
triton_valr=   s         rz   Ú_shaped_constantz TritonOverrides._shaped_constant“  s[   € ä×#Ñ#×1Ñ1°%Ó8ˆÜ"¡5¨£<Ó0ˆ
Ü)¨%Ó0ˆà˜,Ò&àÐð ˜%˜  : ,¨b°°¸QÐ?Ð?r|   c                ó*   — | j                  ||g ¬«      S )N©rë   )rO  )rv   r  rè  s      rz   ÚconstantzTritonOverrides.constant¡  s   € à×#Ñ# E¨5¸Ð#Ó;Ð;r|   c                ó   — d| › d�S )NrÃ  rü   r…   ©r  s    rz   ÚabszTritonOverrides.abs¥  ó   € ð ˜a˜S Ð"Ð"r|   c                óÞ   — d| › d|› d�}t        | «      st        |«      rMt        «       j                  | |«      }|t        j                  t        j
                  fv r|› dt        |«      › d�}|S )Nr®  r�  rü   r  )rþ  r  Útruedivrs   r  rt   r=   ©r  ÚyÚoutr;  s       rz   rX  zTritonOverrides.truediv­  sl   € à�!��C˜�s˜!ˆnˆÜ Ô"Ô&:¸1Ô&=Ü)Ó+×3Ñ3°A°qÓ9ˆIØœUŸ]™]¬E¯M©MÐ:Ñ:Ø˜˜T¤+¨iÓ"8Ð!9¸Ð;�àˆ
r|   c                óÞ   — d| › d|› d�}t        | «      st        |«      rMt        «       j                  | |«      }|t        j                  t        j
                  fv r|› dt        |«      › d�}|S )Nr®  r‚  rü   r  )rþ  r  Úmodrs   r  rt   r=   rY  s       rz   r]  zTritonOverrides.mod·  sl   € à�!��C˜�s˜!ˆnˆÜ Ô"Ô&:¸1Ô&=Ü)Ó+×/Ñ/°°1Ó5ˆIØœUŸ]™]¬E¯M©MÐ:Ñ:Ø˜˜T¤+¨iÓ"8Ð!9¸Ð;�Øˆ
r|   c                ó   — d| › d�S )Nzlibdevice.abs(rü   r…   rT  s    rz   Úlibdevice_abszTritonOverrides.libdevice_absÀ  ó   € ð   ˜s !Ð$Ð$r|   c                ó\   — t         j                  rd| › dt        j                  › d�S d| › d�S )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.
        úlibdevice.exp2(rß  rü   ztl_math.exp()r   Úuse_fast_mathr-  Ú_LOG_2_ErT  s    rz   ÚexpzTritonOverrides.expÅ  s8   € ô ×ÒØ$ Q C s¬?×+CÑ+CÐ*DÀAÐFÐFà! !  AÐ&Ð&r|   c                ó   — d| › d�S )Nzlibdevice.exp(rü   r…   rT  s    rz   Úlibdevice_expzTritonOverrides.libdevice_expÔ  r`  r|   c                ó   — d| › d�S )Nrb  rü   r…   rT  s    rz   Úexp2zTritonOverrides.exp2Ù  ó   € ð !   1Ð%Ð%r|   c                ó   — d| › d�S )Nzlibdevice.expm1(rü   r…   rT  s    rz   Úexpm1zTritonOverrides.expm1Þ  ó   € ð " !  AÐ&Ð&r|   c                ó   — d| › d�S ©Nzlibdevice.sqrt(rü   r…   rT  s    rz   ÚsqrtzTritonOverrides.sqrtã  rj  r|   c                ó   — d| › d�S ro  r…   rT  s    rz   Úlibdevice_sqrtzTritonOverrides.libdevice_sqrtè  rj  r|   c                ó  — t         j                  j                  }|dk(  ry|dk(  r	d| › d| › d�S |dk(  r| › d�S |€8t        j                  t        j
                  d	t        j                  «      | «      S t        d
|›�«      ‚)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ÚmaximumrR  rs   rõ  ÚAssertionError)r  Úbugs     rz   Úreluz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| › d|› d�S )Nztriton_helpers.minimum(rû   rü   r…   ©r¹  rº  s     rz   ÚminimumzTritonOverrides.minimumÿ  ó   € à(¨¨¨2¨a¨S°Ð2Ð2r|   c                ó   — d| › d|› d�S )Nztriton_helpers.maximum(rû   rü   r…   r~  s     rz   ry  zTritonOverrides.maximum  r€  r|   c                ó   — d| › d|› d|› d�S )Nr¦  rû   rü   r…   )r¹  rº  r¨  s      rz   ÚwherezTritonOverrides.where  s   € à˜1˜#˜R ˜s " Q C qÐ)Ð)r|   rE   )Úconstraintsrè  Úis_pureÚpackc                óð   — t        |«      }dj                  |D �cg c]  }t        |«      ‘Œ c}«      }|€#dj                  dg|D �	cg c]  }	d‘Œ c}	z   «      }d| › d|› d|› d|› d|› d	|› d
�S c c}w c c}	w )Nrû   z=rrÔ   ztl.inline_asm_elementwise('z', 'z', [z	], dtype=z
, is_pure=z, pack=rü   )ré  rØ   r´   )
Úasmr„  rè  r…  r†  Úinputsr=   ÚiÚ
input_refsÚ_s
             rz   Úinline_asm_elementwisez&TritonOverrides.inline_asm_elementwise  s¦   € ô *¨%Ó0ˆØ—Y‘Y°Ö7¨1¤ A¥Ò7Ó8ˆ
ØÐØŸ)™) T F¸6Ö-B°aªcÒ-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   ›A.Á	A3
c                ó   — d| › d�S )Nztl_math.cos(rü   r…   rT  s    rz   ÚcoszTritonOverrides.cos  rV  r|   c                ó   — d| › d�S )Nzlibdevice.cos(rü   r…   rT  s    rz   Úlibdevice_coszTritonOverrides.libdevice_cos  r`  r|   c                ó   — d| › d�S )Nztl_math.sin(rü   r…   rT  s    rz   ÚsinzTritonOverrides.sin  rV  r|   c                ó   — d| › d�S )Nzlibdevice.sin(rü   r…   rT  s    rz   Úlibdevice_sinzTritonOverrides.libdevice_sin$  r`  r|   c                ó   — t        d«      ‚)Nz/ops.index_expr not implemented outside a kernel©ÚNotImplementedError)rv   r  rè  s      rz   Ú
index_exprzTritonOverrides.index_expr)  s   € ä!Ð"SÓTÐTr|   c                ó   — t        d«      ‚)Nz+ops.masked not implemented outside a kernelr—  )rÌ   ÚbodyÚothers      rz   ÚmaskedzTritonOverrides.masked-  s   € ä!Ð"OÓPÐPr|   c                ó   — d| › d�S )Nzlibdevice.lgamma(rü   r…   rT  s    rz   ÚlgammazTritonOverrides.lgamma1  ó   € ð # 1 # QÐ'Ð'r|   c                ó   — d| › d�S )Nzlibdevice.erf(rü   r…   rT  s    rz   ÚerfzTritonOverrides.erf6  r`  r|   c                ó   — d| › d�S )Nzlibdevice.cosh(rü   r…   rT  s    rz   ÚcoshzTritonOverrides.cosh;  rj  r|   c                ó   — d| › d�S )Nzlibdevice.sinh(rü   r…   rT  s    rz   ÚsinhzTritonOverrides.sinh@  rj  r|   c                ó   — d| › d�S )Nzlibdevice.acos(rü   r…   rT  s    rz   ÚacoszTritonOverrides.acosE  rj  r|   c                ó   — d| › d�S )Nzlibdevice.acosh(rü   r…   rT  s    rz   ÚacoshzTritonOverrides.acoshJ  rm  r|   c                ó   — d| › d�S )Nzlibdevice.asin(rü   r…   rT  s    rz   ÚasinzTritonOverrides.asinO  rj  r|   c                ó   — d| › d�S )Nzlibdevice.asinh(rü   r…   rT  s    rz   ÚasinhzTritonOverrides.asinhT  rm  r|   c                ó   — d| › d|› d�S )Nzlibdevice.atan2(rû   rü   r…   ©r  rZ  s     rz   Úatan2zTritonOverrides.atan2Y  ó   € ð " !  B q c¨Ð+Ð+r|   c                ó   — d| › d�S )Nzlibdevice.atan(rü   r…   rT  s    rz   ÚatanzTritonOverrides.atan^  rj  r|   c                ó   — d| › d�S )Nzlibdevice.atanh(rü   r…   rT  s    rz   ÚatanhzTritonOverrides.atanhc  rm  r|   c                ó   — d| › d|› d�S )Nzlibdevice.copysign(rû   rü   r…   r°  s     rz   ÚcopysignzTritonOverrides.copysignh  s   € ð % Q C r¨!¨¨AÐ.Ð.r|   c                ó   — d| › d�S )Nzlibdevice.erfc(rü   r…   rT  s    rz   ÚerfczTritonOverrides.erfcm  rj  r|   c                ó   — d| › d�S )Nzlibdevice.erfinv(rü   r…   rT  s    rz   ÚerfinvzTritonOverrides.erfinvr  r   r|   c                ó   — d| › d|› d�S )Nzlibdevice.hypot(rû   rü   r…   r°  s     rz   ÚhypotzTritonOverrides.hypotw  r²  r|   c                ó   — d| › d�S )Nzlibdevice.log10(rü   r…   rT  s    rz   Úlog10zTritonOverrides.log10|  rm  r|   c                ó   — d| › d�S )Nzlibdevice.log2(rü   r…   rT  s    rz   Úlog2zTritonOverrides.log2�  rj  r|   c                ó   — d| › d|› d�S )Nzlibdevice.nextafter(rû   rü   r…   r°  s     rz   Ú	nextafterzTritonOverrides.nextafter†  s   € ð & a S¨¨1¨#¨QÐ/Ð/r|   c                ó   — | › d|› �S ©Nr×   r…   r~  s     rz   Úlogical_andzTritonOverrides.logical_and‹  ó   € à��C˜�sˆ|Ðr|   c                ó   — | › d�S )Nz == 0r…   ©r¹  s    rz   Úlogical_notzTritonOverrides.logical_not�  s   € à��Eˆ{Ðr|   c                ó   — | › d|› �S ©Nz | r…   r~  s     rz   Ú
logical_orzTritonOverrides.logical_or“  rÈ  r|   c                ó   — d| › d|› d�S )Nr®  ú ^ rü   r…   r~  s     rz   Úlogical_xorzTritonOverrides.logical_xor—  s   € à�1�#�S˜˜˜1ˆ~Ðr|   c                ó   — | › d|› �S rÆ  r…   r~  s     rz   Úbitwise_andzTritonOverrides.bitwise_and›  rÈ  r|   c                ó   — d| › �S )Nú~r…   rÊ  s    rz   Úbitwise_notzTritonOverrides.bitwise_notŸ  s   € à�1�#ˆwˆr|   c                ó   — | › d|› �S rÍ  r…   r~  s     rz   Ú
bitwise_orzTritonOverrides.bitwise_or£  rÈ  r|   c                ó   — | › d|› �S )NrÐ  r…   r~  s     rz   Úbitwise_xorzTritonOverrides.bitwise_xor§  rÈ  r|   c                ó   — | › d|› �S )Nz << r…   r~  s     rz   Úbitwise_left_shiftz"TritonOverrides.bitwise_left_shift«  ó   € à��D˜˜ˆ}Ðr|   c                ó   — | › d|› �S )Nz >> r…   r~  s     rz   Úbitwise_right_shiftz#TritonOverrides.bitwise_right_shift¯  rÝ  r|   c                ó    — d|› d�}d| › d|› d�S )Nr®  ú).to(tl.uint32)ztl.rand(rû   rü   r…   ©Úseedr—   s     rz   ÚrandzTritonOverrides.rand³  s%   € à�V�H˜OÐ,ˆØ˜$˜˜r & ¨Ð+Ð+r|   c                ó    — d|› d�}d| › d|› d�S )Nr®  rá  z	tl.randn(rû   rü   r…   râ  s     rz   ÚrandnzTritonOverrides.randn¸  s%   € à�V�H˜OÐ,ˆØ˜4˜&  6 (¨!Ð,Ð,r|   c           	     ó,   — d|› d�}d| › d|› d|› d|› d�	S )Nr®  rá  ztriton_helpers.randint64(rû   rü   r…   )rã  r—   ÚlowÚhighs       rz   Ú	randint64zTritonOverrides.randint64½  s1   € à�V�H˜OÐ,ˆØ*¨4¨&°°6°(¸"¸S¸EÀÀDÀ6ÈÐKÐKr|   c                ó   — t        d«      ‚)Nz.ops.load_seed not implemented outside a kernelr—  )rA  r—   s     rz   Ú	load_seedzTritonOverrides.load_seedÂ  s   € ä!Ð"RÓSÐSr|   c                ó   — d| › d�S )Nzlibdevice.rsqrt(rü   r…   rT  s    rz   ÚrsqrtzTritonOverrides.rsqrtÆ  rm  r|   c                ó   — d| › d�S )Nzlibdevice.log1p(rü   r…   rT  s    rz   Úlog1pzTritonOverrides.log1pË  rm  r|   c                ó   — d| › d�S )Nzlibdevice.tan(rü   r…   rT  s    rz   ÚtanzTritonOverrides.tanÐ  r`  r|   c                ó   — d| › d�S )Nzlibdevice.tanh(rü   r…   rT  s    rz   ÚtanhzTritonOverrides.tanhÕ  rj  r|   c                ó   — d| › d�S )Nztl.sigmoid(rü   r…   rT  s    rz   ÚsigmoidzTritonOverrides.sigmoidÚ  s   € ð ˜Q˜C˜qÐ!Ð!r|   c                ó   — d| › d| › d| › d�S )Nz(libdevice.signbit(z) != 0) if (z).dtype is tl.float32 else z < 0r…   rT  s    rz   ÚsignbitzTritonOverrides.signbitß  s#   € ð " !  L°°Ð3NÈqÈcÐQUÐVð	
r|   c                ó   — d| › d|› d�S )Nzlibdevice.fmod(rû   rü   r…   r~  s     rz   ÚfmodzTritonOverrides.fmodæ  s   € ð !   2 a S¨Ð*Ð*r|   c                ó   — d| › d|› d�S )Nr£  rû   rü   r…   r~  s     rz   ÚpowzTritonOverrides.powë  s   € ð   ˜s " Q C qÐ)Ð)r|   c                ó   — d| › d�S )Nztl_math.log(rü   r…   rT  s    rz   ÚlogzTritonOverrides.logð  rV  r|   c                ó   — d| › d�S )Nzlibdevice.log(rü   r…   rT  s    rz   Úlibdevice_logzTritonOverrides.libdevice_logõ  r`  r|   F)rx   c                ó   — d| › d�S )Nzlibdevice.isinf(ú).to(tl.int1)r…   rT  s    rz   ÚisinfzTritonOverrides.isinfú  ó   € ð " !  MÐ2Ð2r|   c                ó   — d| › d�S )Nzlibdevice.isnan(r  r…   rT  s    rz   ÚisnanzTritonOverrides.isnanÿ  r  r|   c                ó   — d| › d�S )Nzlibdevice.nearbyint(rü   r…   rT  s    rz   ÚroundzTritonOverrides.round  s   € ð & a S¨Ð*Ð*r|   c                ó   — d| › d�S )Nr”  rü   r…   rT  s    rz   ÚfloorzTritonOverrides.floor	  rm  r|   c                óH   — | › d|› �}| › d|› �}d| › d|› d|› d|› d|› d|› d	�S )
NrŒ  r‚  z
tl.where((z
 < 0) != (z < 0), tl.where(z != 0, z - 1, ú), rü   r…   )r¹  rº  r†  Úrems       rz   ÚfloordivzTritonOverrides.floordiv  sV   € ð
 ��D˜˜ˆ}ˆØ��3�q�cˆlˆØ˜A˜3˜j¨¨Ð+;¸C¸5ÀÈÀvÈVÐTXÐSYÐY\Ð]aÐ\bÐbcÐdÐdr|   c                óf  — t        j                  dt        j                  «      }t        j                  t        j
                  || «      t        j                  «      }t        j                  t        j
                  | |«      t        j                  «      }t        j                  ||«      }|› d| › d�S )Nr   r  ú.dtype))rx  rR  rs   rõ  r<  Últrð  Úsub)r  ÚzÚleftÚrightr  s        rz   ÚsignzTritonOverrides.sign  su   € ä�L‰L˜œEŸK™KÓ(ˆÜ�|‰|œSŸV™V A q›\¬E¯J©JÓ7ˆÜ—‘œcŸf™f Q¨›l¬U¯Z©ZÓ8ˆÜ�g‰g�d˜EÓ"ˆØ��d˜1˜#˜WÐ%Ð%r|   c                ó   — d| › d�S )Nrn  rü   r…   rT  s    rz   ÚtrunczTritonOverrides.trunc  rm  r|   c                ó   — | › d|› �S )NrŒ  r…   r~  s     rz   ÚtruncdivzTritonOverrides.truncdiv$  s   € ð ��D˜˜ˆ}Ðr|   c                ó   — d| › d�S )Nrš  rü   r…   rT  s    rz   ÚceilzTritonOverrides.ceil*  rj  r|   )NT)rè  r  r2  zOptional[torch.dtype])rè  r  r2  r  )Xrr   r€   r�   r‚   ÚmathrÂ  Úerd  r\  r<  rH  rO  r„   rR  r+  rU  rX  r]  r_  re  rg  ri  rl  rp  rr  r|  r  ry  rƒ  rs   rt   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  r
  r  r  r  r  r  r…   r|   rz   r-  r-  :  s3  „ Ù(àˆt�y‰y˜Ÿ™Ó €Hàð ,0Øð	6&àð6&ð )ò6&ó ð6&ðp ò@ó ð@ð4 ñ@ó ð@ð ñ<ó ð<ð ÙÓñ#ó ó ð#ð ñó ðð ñó ðð ÙÓñ%ó ó ð%ð ÙÓñ'ó ó ð'ð ÙÓñ%ó ó ð%ð ÙÓñ&ó ó ð&ð ÙÓñ'ó ó ð'ð ÙÓñ&ó ó ð&ð ÙÓñ&ó ó ð&ð ñó ðð" ñ3ó ð3ð ñ3ó ð3ð ñ*ó ð*ð à"&¨e¯m©mÀTÐPQóKó ðKð ÙÓñ#ó ó ð#ð ÙÓñ%ó ó ð%ð ÙÓñ#ó ó ð#ð ÙÓñ%ó ó ð%ð ñUó ðUð ñQó ðQð ÙÓñ(ó ó ð(ð ÙÓñ%ó ó ð%ð ÙÓñ&ó ó ð&ð ÙÓñ&ó ó ð&ð ÙÓñ&ó ó ð&ð ÙÓñ'ó ó ð'ð ÙÓñ&ó ó ð&ð ÙÓñ'ó ó ð'ð ÙÓñ,ó ó ð,ð ÙÓñ&ó ó ð&ð ÙÓñ'ó ó ð'ð ÙÓñ/ó ó ð/ð ÙÓñ&ó ó ð&ð ÙÓñ(ó ó ð(ð ÙÓñ,ó ó ð,ð ÙÓñ'ó ó ð'ð ÙÓñ&ó ó ð&ð ÙÓñ0ó ó ð0ð ñó ðð ñó ðð ñó ðð ñó ðð ñó ðð ñó ðð ñó ðð ñó ðð ñó ðð ñó ðð ñ,ó ð,ð ñ-ó ð-ð ñLó ðLð ñTó ðTð ÙÓñ'ó ó ð'ð ÙÓñ'ó ó ð'ð ÙÓñ%ó ó ð%ð ÙÓñ&ó ó ð&ð ÙÓñ"ó ó ð"ð ñ
ó ð
ð ÙÓñ+ó ó ð+ð ÙÓñ*ó ó ð*ð ÙÓñ#ó ó ð#ð ÙÓñ%ó ó ð%ð Ù¨Ô/ñ3ó 0ó ð3ð Ù¨Ô/ñ3ó 0ó ð3ð ÙÓñ+ó ó ð+ð ÙÓñ'ó ó ð'ð ñeó ðeð ñ&ó ð&ð ÙÓñ'ó ó ð'ð ñó ðð
 ÙÓñ&ó ó ñ&r|   r-  rŒ   c                  ó`   — e Zd ZdZed„ «       Zed„ «       Zed„ «       Zed„ «       Z	ed„ «       Z
y)ÚTritonKernelOverridesa   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                ór   — t         j                  j                  «       }dg|z  }| j                  |||¬«      S )NrE   rQ  )rC   r  Útriton_tensor_ndimrO  )rv   r  rè  Úndimrë   s        rz   rR  zTritonKernelOverrides.constant;  s9   € ô
 �x‰x×*Ñ*Ó,ˆØ��d‘
ˆØ×#Ñ# E¨5¸Ð#Ó>Ð>r|   c                ón  — t         j                  j                  |d¬«      }t        |t        «      sJ ‚t         j                  j
                  dk(  rt        j                  nt        j                  }|t        j                  t        j                  fvr|n|}t        j                  j                  }	 dt        j                  _        t         j                  j                  j                  t         j                  j                  |j                  t!        |«      |¬«      }|t        j                  _        |t        j                  t        j                  fvr^t         j                  j                  j                  t         j                  j                  | j#                  ||«      t%        |«      ¬«      }nÖ|}|j&                  D ]l  }t)        |t*        j,                  «      sŒt        j.                  |t         j                  j                  j0                  |j2                     j4                  «      }Œn ||k7  rTt         j                  j                  j                  t         j                  j                  | j#                  ||«      |¬«      }|j6                  |_        |S # |t        j                  _        w xY w)NF©Ú	block_ptrútl.int32©r  rè  rç  )rC   r  Úindexingrb  r³   rr  rs   rõ  Úint64r   Útest_configsÚruntime_triton_dtype_assertÚcseÚgenerateÚcomputerµ   r4   r<  r?   Úfree_symbolsr   r   rÂ   Úpromote_typesÚvarname_maprA  rè  r·   )rv   r  rè  r)  rr  Úorigrý  Ú	index_vars           rz   r™  z TritonKernelOverrides.index_exprD  sû  € ä—8‘8×$Ñ$ T°UÐ$Ó;ˆÜ˜(¤OÔ4Ð4Ð4ô &'§X¡X×%9Ñ%9¸ZÒ%G”e—k’kÌUÏ[É[ˆØ¤u§{¡{´E·K±KÐ&@Ñ@‘Àkˆô ×"Ñ"×>Ñ>ˆð		CØ>CŒF×ÑÔ;Ü—(‘(—,‘,×'Ñ'Ü—‘× Ñ Ø×"Ñ"Ü,¨TÓ2Øð	 (ó ˆCð ?CŒF×ÑÔ;àœŸ™¤e§k¡kÐ2Ñ2Ü—(‘(—,‘,×'Ñ'Ü—‘× Ñ Ø—‘˜S %Ó(Ü)¨%Ó0ð (ó ‰Cð  ˆEØ!×.Ñ.ò �	Ü! )¬T¯X©XÕ6Ü!×/Ñ/ØœqŸx™xŸ|™|×7Ñ7¸	¿¹ÑG×MÑMó‘Eðð ˜Ò#Ü—h‘h—l‘l×+Ñ+Ü—H‘H×$Ñ$Ø—L‘L  kÓ2Ø%ð ,ó �ð !×*Ñ*ˆŒØˆ
øð9 ?CŒF×ÑÕ;ús   Â2A-J ÊJ4c           
     óÜ  — | �ot         j                  j                  �Ut        j                  j
                  j                  t        j                  j                  | › d�t         j                  ¬«      } |j                  j                  d¬«      }|sJ d«       ‚d}|D ]>  }|j                  D ]-  }|j                  dk7  st        |j                  d   «      sŒ+d	} Œ> Œ@ |rd n|}t        j                  j                  | |¬
«      5 } |«       }	d d d «       |r®	j                  j                   rt        |«      }t        j                  j
                  j                  t        j                  j                  d|	› dt#        |«      › d|	› d�t%        j&                  |«      |	j(                  ¬«      }t+        j,                  |	|«      }
n	}
|
j.                  j1                  «       |
S # 1 sw Y   ŒØxY w)Nú.to(tl.int1)rç  Úoutput)Úopz)graph for body does not contain an outputFÚloadrE   T©r  rJ  z.shape, rû   r  r(  )rs   rx  ry  rC   r  r-  r.  r/  r}   r  Ú
find_nodesrC  Útargetr_   Ú
mask_loadsr  Úis_boolrV   r   Úwraprè  rx  rƒ  r·   Údiscard)rÌ   r›  rœ  ÚnodesÚ
need_whereÚnoder	  r  Únew_maskr5  rz  s              rz   r�  zTritonKernelOverrides.maskedw  s®  € àÐ¤§¡× 1Ñ 1Ð =Ü—8‘8—<‘<×(Ñ(Ü—‘× Ñ Ø�&˜Ð%Ü—j‘jð )ó ˆDð —
‘
×%Ñ%¨Ð%Ó2ˆÙÐAÐAÓAˆuàˆ
ð ò 	ˆDØ—y‘yò �Ø—:‘: Ò'Ô+CÀCÇHÁHÈQÁKÕ+PØ!%�JÙñð	ñ #‘¨ˆä�X‰X× Ñ  ¨UÐ Ó3ð 	°xÙ“VˆF÷	ñ à�}‰}×$Ò$Ü˜U›�ä—H‘H—L‘L×)Ñ)Ü—‘× Ñ Ø˜6˜( (¬=¸Ó+?Ð*@ÀÀ6À(È'ÐRÜ"×'Ñ'¨Ó.Ø—l‘lð	 *ó ˆEô —)‘)˜H f¨eÓ4‰CàˆCà�‰×Ñ˜hÔ'Øˆ
÷'	ð 	ús   ÄG"Ç"G+c                ó¶   — t         j                  j                  j                  | «      }d|› dt         j                  j                  j	                  d|«      › d�S )Nútl.load(ú + Úload_seed_offsetrü   )rC   r  rC  ÚinputÚseed_offset)rA  r—   rý  s      rz   rì  zTritonKernelOverrides.load_seed¤  sI   € ä�h‰h�m‰m×!Ñ! $Ó'ˆà�s�e˜3œqŸx™xŸ}™}×8Ñ8Ð9KÈVÓTÐUÐUVÐWð	
r|   c                ó  — d| › d�}t         j                  j                  j                  |«      x}r|S t         j                  j                  j	                  | j
                  ¬«      }t         j                  j                  j	                  t        j                  ¬«      }t         j                  j                  j                  |› d|› d| › d�«       t         j                  j                  j                  |||f«       ||fS )Nzfrexp(rü   rç  rû   z = triton_helpers.frexp()rC   r  r-  Útry_getÚnewvarrè  rs   rõ  r/  r�   Úput)r  Ú	cache_keyÚcse_valÚmantissaÚexponents        rz   ÚfrexpzTritonKernelOverrides.frexp«  sÉ   € à˜Q˜C˜q�Mˆ	Ü—h‘h—l‘l×*Ñ*¨9Ó5Ð5ˆ7Ð5ØˆNä—8‘8—<‘<×&Ñ&¨Q¯W©WÐ&Ó5ˆÜ—8‘8—<‘<×&Ñ&¬U¯[©[Ð&Ó9ˆÜ	�‰×Ñ×"Ñ"Øˆj˜˜8˜*Ð$<¸Q¸C¸qÐAô	
ô 	
�‰�‰×Ñ˜ X¨xÐ$8Ô9Ø˜(Ð#Ð#r|   N)rr   r€   r�   r‚   r„   rR  r™  r\  r�  rì  rS  r…   r|   rz   r   r   3  sm   „ ñð ñ?ó ð?ð ñ0ó ð0ðd ñ*ó ð*ðX ñ
ó ð
ð ñ$ó ñ$r|   r   c                  óH   — e Zd ZU dZded<   ded<   dd„Zddœdd	„Zd
„ Zd„ Zy)ÚHelperFunctionsz#An ordered set of helper functions.zdict[str, str]Ú_templates_seenú	list[str]Úfinalized_helpersc                ó    — i | _         g | _        y rq   )rV  rX  r¾   s    rz   r  zHelperFunctions.__init__À  s   € Ø!ˆÔØ!#ˆÕr|   Ú_triton_helper_fn©Ú	base_namec               óê   — | j                   j                  |«      }|�|S |› t        | j                  «      › �}|| j                   |<   | j                  j	                  |j                  |¬«      «       |S )a9  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.

        )rA  )rV  Úgetr  rX  rc  rD  )r¿   Útemplate_coder\  Úexisting_namerA  s        rz   ÚaddzHelperFunctions.addÄ  sw   € ð ×,Ñ,×0Ñ0°Ó?ˆØÐ$à Ð à�œS ×!7Ñ!7Ó8Ð9Ð:ˆØ.2ˆ×Ñ˜]Ñ+Ø×Ñ×%Ñ% m×&:Ñ&:ÀÐ&:Ó&EÔFØˆr|   c                ó,   — t        | j                  «      S rq   )ÚiterrX  r¾   s    rz   Ú__iter__zHelperFunctions.__iter__Ú  s   € Ü�D×*Ñ*Ó+Ð+r|   c                ó    — | j                   |   S rq   )rX  )r¿   rJ  s     rz   Ú__getitem__zHelperFunctions.__getitem__Ý  s   € Ø×%Ñ% cÑ*Ð*r|   N©r~   r   )r_  r´   r~   r´   )	rr   r€   r�   r‚   rƒ   r  ra  rd  rf  r…   r|   rz   rU  rU  º  s+   … Ù-à#Ó#Ø Ó ó$ð 4Gõ ò,,ó+r|   rU  c                  óÚ   — e Zd ZU dZ ej
                  e¬«      Zded<    ej
                  e¬«      Z	ded<    ej
                  e¬«      Z
ded<    ej
                  e¬«      Zded<   d
d„Zy	)rà   zM
    Class representing ND block dimensions, for block pointer analysis.
    )Údefault_factoryr  rë   rí   rï   rò   c                ó�   — t        | «      }t        d„ | |fD «       «      \  }} |di |D �ci c]  }|||   ||   z   “Œ c}¤ŽS c c}w )z0
        Concatenates block parameters.
        c              3  óF   K  — | ]  }t        j                  |«      –— Œ y ­wrq   )r  r  r¯  s     rz   rÍ   z*BlockParameters.__add__.<locals>.<genexpr>ñ  s   è ø€ ÒB¨q”[×'Ñ'¨×*ÑBùs   ‚!r…   )rµ  r¶  )r¿   rœ  rv   r¹  rº  r2  s         rz   Ú__add__zBlockParameters.__add__ì  sR   € ô �4‹jˆÜÑB°T¸5°MÔBÓB‰ˆˆ1ÙÑ9°aÖ8¨s�c˜1˜S™6 A c¡F™?Ñ*Ò8Ñ9Ð9ùÒ8s   «AN)rœ  rà   r~   rà   )rr   r€   r�   r‚   r  Úfieldr(  rë   rƒ   rí   rï   rò   rl  r…   r|   rz   rà   rà   á  sn   … ñð 0˜k×/Ñ/ÀÔE€EÐÓEØ$5 K×$5Ñ$5ÀdÔ$K€KÐ!ÓKØ 1 × 1Ñ 1À$Ô G€GÐÓGØ 1 × 1Ñ 1À$Ô G€GÐÓGô:r|   rà   c                  ó*   — e Zd ZdZd„ Zdd„Zd„ Zd„ Zy)Ú"CooperativeReductionWorkspaceCachez‘
    The scratch space used for cooperative reductions can be reused
    after two reduction loops.  This keeps track of what can be reused.
    c                óš   — || _         g | _        g | _        t        j                  t        j
                  «      | _        d| _        d| _        y r=  )	rC  Úcurrent_loopÚ
prior_loopÚcollectionsÚdefaultdictÚdequeÚready_for_reuseÚ
loop_countÚstore_count)r¿   rC  s     rz   r  z+CooperativeReductionWorkspaceCache.__init__û  s@   € ØˆŒ	ØˆÔØˆŒÜ*×6Ñ6´{×7HÑ7HÓIˆÔØˆŒØˆÕr|   c                óÞ   — | j                   j                  |«      }|r|j                  «       S | j                  j	                  |d«      \  }}| j
                  j                  |||f«       ||fS rR  )rv  r^  ÚpopleftrC  Ú	workspacerq  rc  )r¿   ÚnbytesÚcachedÚws_nameÚ	ws_offsets        rz   Úallocatez+CooperativeReductionWorkspaceCache.allocate  si   € Ø×%Ñ%×)Ñ)¨&Ó1ˆÙØ—>‘>Ó#Ð#Ø!ŸY™Y×0Ñ0°¸Ó?Ñˆ�Ø×Ñ× Ñ  &¨'°9Ð!=Ô>Ø˜Ð#Ð#r|   c                óÈ   — | j                   D ]&  \  }}}| j                  |   j                  ||f«       Œ( | j                  | _         g | _        | xj                  dz  c_        y ©NrE   )rr  rv  rc  rq  rw  )r¿   r|  r~  r  s       rz   Úon_loop_endz.CooperativeReductionWorkspaceCache.on_loop_end  s_   € à*.¯/©/ò 	FÑ&ˆF�G˜YØ× Ñ  Ñ(×/Ñ/°¸)Ð0DÕEð	Fà×+Ñ+ˆŒØˆÔØ�Š˜1ÑŽr|   c                óH   — | j                   }| xj                   dz  c_         |S r‚  )rx  )r¿   Úpriors     rz   Úincrement_store_countz8CooperativeReductionWorkspaceCache.increment_store_count  s#   € Ø× Ñ ˆØ×Ò˜AÑÕØˆr|   N)r|  r»   )rr   r€   r�   r‚   r  r€  rƒ  r†  r…   r|   rz   ro  ro  õ  s   „ ñò
ó$òór|   ro  c                  ó$   — e Zd ZU ded<   d„ Zd„ Zy)ÚFixedTritonConfigzdict[str, int]r   c                ó    — | j                   |   S rq   ©r   ©r¿   r  s     rz   rf  zFixedTritonConfig.__getitem__  s   € Ø�{‰{˜4Ñ Ð r|   c                ó   — || j                   v S rq   rŠ  r‹  s     rz   Ú__contains__zFixedTritonConfig.__contains__   s   € Ø�t—{‘{Ð"Ð"r|   N)rr   r€   r�   rƒ   rf  r�  r…   r|   rz   rˆ  rˆ    s   … àÓò!ó#r|   rˆ  c                  ó   — e Zd ZdZdd„Zy)Ú	TritonCSEz…
    Subclasses CSE to apply the current load mask to the cache key to avoid CSEing
    variables across separate masked blocks.
    c                óZ   — t         j                  j                  x}r||j                  fS |S rq   )rC   r  Ú
_load_maskrA  )r¿   rO  rÌ   s      rz   Úaugment_keyzTritonCSE.augment_key*  s,   € Ü—8‘8×&Ñ&Ð&ˆ4Ð&Ø˜tŸy™yÐ)Ð)àÐr|   N)rO  r´   r~   zUnion[str, tuple[str, str]])rr   r€   r�   r‚   r’  r…   r|   rz   r�  r�  $  s   „ ñô
r|   r�  c                  ó*  ‡ — e Zd ZU eZded<   eZded<   dZ	 	 	 dI	 	 	 	 	 dJˆ fd„Z	dKd„Z
dLd	„Zd
„ Zd„ Zd„ Zd„ ZdLd„Zd„ ZedMd„«       Zdddddœ	 dNd„Z	 dO	 	 	 	 	 	 	 dPd„ZdOd„Z	 	 	 	 	 	 	 	 dQd„Zd„ ZdRd„Z	 dS	 	 	 	 	 	 	 	 	 dTd„Zd„ Z	 	 dU	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 dVd„ZdMd„ZdWd„Z	 	 	 	 	 	 	 	 	 	 dXd„Z	 	 dYd„Z dYd „Z!d!„ Z"d"„ Z#d#„ Z$d$„ Z%d%„ Z&	 	 	 	 	 	 dZd&„Z'd[d'„Z(	 	 	 	 	 	 	 	 d\d(„Z)	 	 	 	 	 	 	 	 	 	 d]d)„Z*d*„ Z+d^d+„Z,d,„ Z-d-„ Z.d.„ Z/e0d/„ «       Z1dSd0„Z2e0d1„ «       Z3e0d2„ «       Z4d3„ Z5d_d4„Z6d5„ Z7dSd`d6„Z8dad7„Z9dbd8„Z:dcd9„Z;ddd:„Z<	 	 	 	 	 	 ded;„Z=ddd<„Z>dfd=„Z?dgd>„Z@dhd?„ZAdLd@„ZBdidA„ZCeDd^dB„«       ZEdjdC„ZFdkdD„ZGeDdldE„«       ZHdmdF„ZIdjdG„ZJ	 	 	 	 	 	 dndH„ZKˆ xZLS )oÚTritonKernelrU  Úhelper_functionszCallable[[sympy.Expr], str]ÚkexprTNc                ó  •— || _         || _        t        ‰| �  |fi |¤Ž t	        | j
                  | j                  «      | _        t        «       | _	        t        «       | _
        t        t           «       | _        || _        t        j                   «       | _        t%        t&        t&        f   «       | _        t+        «       | _        t/        j0                  t$        «      | _        t/        j4                  «       | _        t        t8           «       | _        d | _        | j>                  r| jA                  | jB                  «       | jD                  r| jG                  «        | jI                  «        | jD                  r| jK                  «        y y rq   )&Úoptimize_maskÚfixed_configr  r  r�  Únewvar_prefixÚsuffixr-  rM   Úpost_loop_combineÚpost_loop_storer   r   Úoutside_loop_varsr8  r  ÚcountÚblock_ptr_idÚdictr´   Úblock_ptr_to_bufferrU  r•  rs  rt  Úpointer_advancementsÚCounterÚ_load_countsr(   Úautotune_hintsÚtriton_metar$  Úcodegen_reduction_numelsr›  Úcooperative_reductionÚinit_cooperative_reductionÚcodegen_range_treeÚinit_cooperative_reduction_mask)r¿   Útilingr8  r˜  r™  r  r  s         €rz   r  zTritonKernel.__init__7  s3  ø€ ð $1ˆÔØ(ˆÔÜ‰Ñ˜Ñ* 6Ò*Ü˜T×/Ñ/°·±Ó=ˆŒÜ1?Ó1AˆÔÜ/=Ó/?ˆÔÜ!+¬C¡Ó!2ˆÔØ#6ˆÔ Ü%ŸO™OÓ-ˆÔÜ#'¬¬S¨¡>Ó#3ˆÔ Ü /Ó 1ˆÔä×#Ñ#¤DÓ)ð 	Ô!ô 7B×6IÑ6IÓ6KˆÔô )¬Ñ6Ó8ˆÔØ59ˆÔà× Ò Ø×)Ñ)¨$¯)©)Ô4à×%Ò%Ø×+Ñ+Ô-à×ÑÔ!à×%Ò%Ø×0Ñ0Õ2ð &r|   c                ó   — t        |«      S rq   )r=   )r¿   rè  s     rz   Údtype_to_strzTritonKernel.dtype_to_str^  s   € Ü˜5Ó!Ð!r|   c                óp   — | j                   xr) t        j                  j                  | j                  «      S rq   )r$  rC   ÚchoicesÚ should_use_cooperative_reductionr&  r¾   s    rz   r²  z-TritonKernel.should_use_cooperative_reductiona  s-   € Ø×$Ñ$ò 
¬¯©×)SÑ)SØ�M‰Mó*
ð 	
r|   c                ó  ‡ — ‰ j                   sJ ‚‰ j                  D ]$  }|j                  €Œ|xj                  dz  c_        Œ& ‰ j                  d   }‰ j                  rt        |‰ j                  d   «      }‰ j                  j                  |«      ‰ _        t        ‰ j                  «      ‰ _
        ‰ j                  j                  d«       t        ˆ fd„‰ j                  D «       «      r‰ j                  j                  d«       yy)z/One time setup code for cooperative reductions.NrE   r  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  óZ   •K  — | ]"  }|j                   r‰j                  |«       –— Œ$ y ­wrq   )r'  Ú_has_constant_mask)r±   r¢   r¿   s     €rz   rÍ   z:TritonKernel.init_cooperative_reduction.<locals>.<genexpr>‚  s0   øè ø€ ò 
àØ× Ò ð ×'Ñ'¨Ó-Ô-ñ
ùs   ƒ(+z>rsplit_end = tl.where(rsplit_end < rnumel, rsplit_end, rnumel))r©  r,  Úgrid_dimr%  r™  r   rC  Ú
semaphoresÚsemaphores_namero  Ú%cooperative_reduction_workspace_cacher›  r�   rÐ   r�   )r¿   r¢   Ú	sem_counts   `  rz   rª  z'TritonKernel.init_cooperative_reductionf  sî   ø€ à×)Ò)Ð)Ð)ð ×$Ñ$ò 	#ˆDØ�}‰}Ñ(Ø—’ Ñ"–ð	#ð —K‘K Ñ$ˆ	Ø×ÒÜ 	¨4×+<Ñ+<¸XÑ+FÓGˆIØ#Ÿy™y×3Ñ3°IÓ>ˆÔÜ5WØ�I‰Ió6
ˆÔ2ð 	�	‰	×Ñð	ô	
ô ó 
à×(Ñ(ô
ô 
ð
 �I‰I×ÑØPõð
r|   c                ó  — d}| j                   s|› d�}| j                  j                  d|› �«       | j                  «       r| j                  j	                  d«       y | j                   rJ ‚| j                  j                  d«       y )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))r   r›  r�   Ú_has_constant_xmaskr�   )r¿   Úrsplit_aranges     rz   r¬  z,TritonKernel.init_cooperative_reduction_mask‹  s{   € Ø>ˆØ�}Š}Ø,˜o¨YÐ7ˆMØ�	‰	×ÑÐ.¨}¨oÐ>Ô?à×#Ñ#Ô%Ø�I‰I×Ñðõð —}’}Ð$Ð$Ø�I‰I×ÑØeõr|   c                ó2  — | j                   D ]q  }|j                  s| j                  || j                  «       Œ,| j                  sŒ9| j                  j                  |j                  › d| j                  |«      › �«       Œs | j                  r‹t        d„ | j                   D «       «      rS| j                  ddd¬«      }| j                  |«      }| j                  j                  d| j                  |«      › �«       y | j                  | j                  «       y y )Nzbase = c              3  ó4   K  — | ]  }|j                   –— Œ y ­wrq   )Úis_loop©r±   r¢   s     rz   rÍ   z2TritonKernel.codegen_range_tree.<locals>.<genexpr>­  s   è ø€ Ò= D�4—<•<Ñ=ùs   ‚ÚbaseTr˜   zrbase = )r,  rÀ  Úiteration_ranges_codegen_headerr›  r$  r�   r!  Úiteration_ranges_ranges_coderÐ   Ú_get_reduction_symbolsÚ_flatten_reduction_indicesr�   r  Úcodegen_reduction_indices)r¿   r¢   Úrn_basesÚrbases       rz   r«  zTritonKernel.codegen_range_tree   sø   € Ø×$Ñ$ò 		ˆDà—<’<Ø×4Ñ4°T¸4¿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                 ó   — y)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    rz   Úneed_numel_argszTritonKernel.need_numel_args¸  s   € ð r|   c                ó†   — | j                   xr4 t        j                  j                  | j                  | j
                  «      S rq   )r$  rC   r±  Úshould_use_persistent_reductionr&  r©  r¾   s    rz   rÍ  z,TritonKernel.should_use_persistent_reductionÂ  s5   € Ø×$Ñ$ò 
¬¯©×)RÑ)RØ�M‰M˜4×5Ñ5ó*
ð 	
r|   c                óô   — | j                   rlt        | j                  «      | j                  dz   k(  rG| j                  r| j                  d   dk(  S t
        j                  j                  | j                  «      S y)NrE   rª   F)	Úpersistent_reductionr  r%  r#  r™  rC   r±  Úwant_no_x_dimr&  r¾   s    rz   rÐ  zTritonKernel.want_no_x_dimÇ  sb   € à×%Ò%Ü�D—K‘KÓ  D×$;Ñ$;¸aÑ$?Ò?à× Ò Ø×(Ñ(¨Ñ2°aÑ7Ð7Ü—9‘9×*Ñ*¨4¯=©=Ó9Ð9Ør|   c                 ó   — y)Nztl.device_assertr…   r¾   s    rz   Úassert_functionzTritonKernel.assert_functionÑ  s   € à!r|   F)Ú
copy_shapeÚdense_indexingÚoverride_maskr&  c          
     óT  ‡ ‡‡‡‡‡— ‰ j                  ‰«      Š‰j                  }d}t        t           «       Š|D �]j  }t	        |t
        j                  «      sJ ‚|xs t        |t        j                  «      }|rŒAt        |t        j                  «      r?‰ j                  j                  |j                     }	‰j                  |	j                   «       Œšt        |t        j"                  t        j$                  t        j&                  t        j(                  t        j*                  t        j,                  f«      r�Œt        j.                  D �
cg c]  }
t        ||
«      r	t0        |
   ‘Œ }}
t3        |«      dk(  sJ d|j                  › �«       ‚‰j5                  |d   › d�«       �Œm t6        j8                  j:                  xs |xs ‰ j<                  duxr ‰dk7  }d}d}t        t           «       }‰ j?                  «       D ]@  }|jA                  |jB                  «      rd}nd}|j5                  |jD                  › d�«       ŒB |rž‰ jF                  r’t6        j8                  jH                  rx|sv‰ j<                  sjt3        ‰|z
  «      dk(  rY‰ jK                  ‰«      sH|rF‰ jL                  dk(  r7	 	 	 	 	 	 dd	„Š	 	 	 	 	 	 dˆ fd
„Š	 	 	 	 	 	 dˆˆfd„Šdˆˆˆˆ fd„} |«       }|�|S d}‰ jO                  ‰«      }t	        ‰t
        jP                  «      r‰|r|› d�n‰ jS                  «       }d|› d|› d�}‰ jT                  r‰ jW                  «       st        dg«      Šn
t        «       Š‰ j<                  r‰j5                  ‰ j<                  «       tY        |‰||‰«      S |r%|s#|r|› d�n‰ jS                  «       }d|› d|› d�}|Šn|s|rd|› d|› d�}|Š|rt        |g«      Š‰ j<                  r‰j5                  ‰ j<                  «       ‰ j[                  ‰«       tY        |‰||‰«      S c c}
w )zO
        Compute the index and mask to pass to tl.load() or tl.store()
        FrE   zAmbiguous type: r   rÌ   NTr'  c                óÖ   — t        j                  | |j                  «       «      }|€yt        |j                  gt
        j                  |«      g|gt
        j                  |«      g¬«      S )z§
                Matches expressions of the form:
                    idx = s * xindex

                This implies stride (s,), and shape (XBLOCK,).
                N©rë   rí   rï   rò   )rF   Úmatch_affine_block_exprÚsymbolrà   Únumelr–   r£   r¦   )r¼   Ú
range_treer.  s      rz   Úmatch_affine_blockz1TritonKernel.indexing.<locals>.match_affine_block%  sl   € ô -×DÑDØ˜:×,Ñ,Ó.ó�ð �>Øä&Ø%×+Ñ+Ð,Ü!.×!=Ñ!=¸jÓ!IÐ JØ#˜HÜ*×;Ñ;¸JÓGÐHô	ð r|   c                óÈ  •‡‡— |j                  «       }t        j                  dt        j                  t        j
                  |g¬«      ¬«      \  }}t        dt        ‰j                  «      | j                  t        ||«      «      | j                  t        |||«      «      z   «      }t        j                  | ||j                  |«      }|€y|\  }}}	t        j                  |«      }
t         j"                  j$                  Š‰j'                  |j(                  «      Št+        ˆˆfd„|
D «       «      ryt,        j/                  |«      }t1        ||
d   «      gt3        |
dd |dd «      D ��cg c]%  \  }}t        j4                  t1        ||«      |«      ‘Œ' c}}z   }|	D �cg c]#  }t7        ||t,        j9                  |«      i«      ‘Œ% }}t;        ||||¬	«      S c c}}w c c}w )
aÍ  
                Matches higher-dimensional blocks coming from FloorDiv and ModularIndexing.

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

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

                Note that dN does not appear in the expression, but we solve for it
                using range tree numels and the other dims.
                zdenom modulo)Úexclude)rv   r   Nc              3  ól   •K  — | ]+  }‰j                  |‰«       xr ‰j                  |«       –— Œ- y ­wrq   )rG  Ústatically_known_power_of_2)r±   rÛ  Ú	max_blockrù   s     €€rz   rÍ   zETritonKernel.indexing.<locals>.match_mod_div_block.<locals>.<genexpr>z  sH   øè ø€ ò ð ð !×=Ñ=¸eÀYÓOÐOò HØ$×@Ñ@ÀÓGÐGóHñùs   ƒ14r   rE   rØ  )rÚ  r®   ÚsymbolsÚ	functoolsÚpartialÚWildr7  r  Úrange_tree_nodesrŸ  r   r   rF   Úmatch_mod_div_block_exprrÛ  Úget_slice_numelsrC   r  rù   râ  r!  rÐ   r–   r£   r   rý   ÚMinr<   r¦   rà   )r¼   rÜ  r4  ÚdenomÚmoduloÚnum_dimsÚmatch_resultr1  rï   Úblock_index_exprsÚslice_numelsÚlinear_block_sizerÛ  r  rí   r  r¥   râ  rù   r¿   s                    @@€rz   Úmatch_mod_div_blockz2TritonKernel.indexing.<locals>.match_mod_div_block;  sì  ú€ ð( '×-Ñ-Ó/�	ô !&§¡Ø"Ü!×)Ñ)¬%¯*©*¸y¸kÔJô!‘��vô ØÜ˜×-Ñ-Ó.àŸ™¤H¨Y¸Ó$>Ó?ØŸ+™+¤o°iÀÈÓ&OÓPñQó	�ô  3×KÑKØ˜9 j×&6Ñ&6¸ó �ð  Ð'Øð !ñ	ØØØ%ä2×CÑCÀDÓI�ô Ÿ7™7×+Ñ+�Ø ŸN™N¨:×+<Ñ+<Ó=�	Üô ð ".ôô ð
  ô %2×$@Ñ$@ÀÓ$LÐ!äÐ-¨|¸A©Ó?ð1ô '*¨,°q°rÐ*:¸DÀÀ¸HÓ&E÷á"˜˜sô —I‘IœgÐ&7¸Ó?ÀÕEóñ1�ð !2ö	3ð ô Ø˜y¬-×*HÑ*HÈÓ*TÐUõð3�ð 3ô 'ØØ +Ø#Ø)ô	ð ùóùò3s   Å,*GÆ (Gc                ó6   •— ‰‰fD ]  } || |«      }|€Œ|c S  y)ze
                Match a block indexing subexpression involving a single range tree.
                Nr…   )r  rÜ  Ú
match_funcÚmatchrÝ  rò  s       €€rz   Úmatch_block_pointer_subexprz:TritonKernel.indexing.<locals>.match_block_pointer_subexpr�  s:   ø€ ð 'Ø'ð#ò %�Jñ ' t¨ZÓ8�EØÑ(Ø$šð%ð r|   c            	     ó€  •— t        ‰‰j                  j                  «       D � �ci c]  \  } }| |j                  “Œ c}} «      }‰j	                  d¬«      }|D �cg c]&  }t        j                  ||j                  «       «      ‘Œ( }}t        d„ |D «       «      }t        «       }t        ||«      D ]@  \  }}t        |j                  |j                  «      «      dkD  r y  ‰||«      }	|	€ y ||	z  }ŒB |t        |«      z
  }
‰j                  ‰«       t         j#                  ||
|‰‰j$                  ¬«      S c c}} w c c}w )NT)Úreorderc              3  ó<   K  — | ]  }|j                  «       –— Œ y ­wrq   )rÚ  rÁ  s     rz   rÍ   zETritonKernel.indexing.<locals>.match_block_pointer.<locals>.<genexpr>¾  s   è ø€ Ò*Q¸T¨4¯;©;¯=Ñ*Qùó   ‚rE   )rá   râ   r,  r·   r-  )r<   rç  r  r  Úactive_range_treesrF   Úget_subexpr_involving_symbolrÚ  r   rà   rý   r  Úintersectionr0  ÚsumÚfilter_masksrß   r6  râ  )ÚvÚtÚindex_relative_to_xyr_indexr,  r¢   Úindex_subexprsÚrange_symbolsÚblock_paramsÚsubexprrá   r—   r¼   r·   rö  r¿   s              €€€€rz   Úmatch_block_pointerz2TritonKernel.indexing.<locals>.match_block_pointer­  sT  ø€ Ü.8Ø°$×2GÑ2G×2MÑ2MÓ2O×P©$¨!¨Q˜A˜qŸv™v™IÓPó/Ð+ð #×5Ñ5¸dÐ5ÓC�ð !,ö	"ð ô (×DÑDØ3°T·[±[³]õð"�ð "ô !+Ñ*QÀ[Ô*QÓ Q�Ü.Ó0�Ü%(¨°nÓ%Eò 
+‘M�D˜'ô ˜=×5Ñ5°g×6JÑ6JÓKÓLÈqÒPÙ#ñ 9¸À$ÓG�FØ�~Ù#Ø  FÑ*‘Lð
+ð 5´s¸>Ó7JÑJ�ð ×!Ñ! )Ô,Ü&×-Ñ-Ø'Ø$*Ø +Ø'Ø"&§.¡.ð .ó ð ùóE Qùò"s   ¥D5Á+D;z.shaperJ  rû   z, tl.int32)Úxmaskrú   rü   ú.shape))r¼   r»   rÜ  rY   r~   úOptional[BlockParameters])r  r»   rÜ  rY   r~   r
  )r~   zOptional[BlockPtrOptions]).Úprepare_indexingr0  r   r´   rb  r®   r¯   r   r–   r©   r   rÂ   r-  r2  rA  r  r·   ÚUNBACKED_INTÚSIZEÚPRECOMPUTED_SIZEÚINDEXÚFLOATÚUNBACKED_FLOATr­   r   r  ra  r   rŒ   rÔ  r‘  rû  rý  Úvar_listr!  Úallow_block_ptrÚuse_block_ptrÚis_indirect_indexingrr  r  r>  Údense_size_strr™  r¼  r³   rÿ  )r¿   r¼   rÓ  rÔ  rÕ  r&  Ú
index_varsrÅ   rý  Úcse_varr    Úprefix_matchesÚ
need_denseÚ
have_denseÚhave_loop_varsÚdense_mask_varsr¢   r  Úoptionsr¹   rµ   r·   rÝ  rö  rò  s   ``                   @@@@rz   r)  zTritonKernel.indexingÕ  s!  ý€ ð ×%Ñ% eÓ,ˆØ×'Ñ'ˆ
Øˆ
äœs‘OÓ%ˆ	Øó 	:ˆCÜ˜c¤5§<¡<Ô0Ð0Ð0Ø#ò ¤~Ø”]×2Ñ2ó(ˆJñ ØÜ ¤T§X¡XÔ.àŸ(™(×.Ñ.¨s¯x©xÑ8�Ø× Ñ  ×!2Ñ!2Õ3ÜØä×%Ñ%Ü—I‘IÜ×)Ñ)Ü—J‘JÜ—J‘JÜ×'Ñ'ðô
ñ ô
 !.× 9Ñ 9ö"àÜ% c¨4Ô0ô ˜tÓ$ð"�ð "ô
 ˜>Ó*¨aÒ/ÐNÐ3CÀCÇHÁHÀ:Ð1NÓNÐ/Ø—‘ °Ñ!2Ð 3°4Ð8Ö9ð?	:ôD �M‰M×(Ñ(ò +Øò+à�‰ dÐ*òð �q‰jð	 	ð ˆ
ØˆÜ$¤S™/Ó+ˆà×+Ñ+Ó-ò 	6ˆDØ×&Ñ& t§}¡}Ô5Ø!%‘à"�
Ø×Ñ 4§;¡; -¨tÐ 4Õ5ð	6ñ Ø×$Ò$Ü—‘×+Ò+Ù!Ø—O’OÜ�I Ñ/Ó0°AÒ5Ø×-Ñ-¨eÔ4Ùà× Ñ  JÒ.ðØ!ðØ/Bðà*óð,`Ø!ð`Ø/Bð`à*õ`ðDØ ðØ.Aðà*ö÷ *ð *ñZ *Ó+ˆGØÐ"Ø�àˆ
Ø×%Ñ% eÓ,ˆ	Ü�eœUŸ]™]Ô+Ù2<˜J˜< vÑ.À$×BUÑBUÓBWˆJØ" : ,¨b°°¸;ÐGˆIØ× Ò ¨×)AÑ)AÔ)CÜ&¨ yÓ1‘	ä&›L�	Ø�ŠØ—‘˜dŸo™oÔ.Ü" 9¨i¸ÀZÐQVÓWÐWá™jÙ2<˜J˜< vÑ.À$×BUÑBUÓBWˆJØ*¨9¨+°R¸
°|À1ÐEˆIØ'‰IÙ¡JØ*¨9¨+°R¸
°|À7ÐKˆIØ'ˆIáÜ" M ?Ó3ˆIà�?Š?Ø�M‰M˜$Ÿ/™/Ô*à×Ñ˜)Ô$ä˜y¨)°ZÀÈUÓSÐSùòy"s   ÅP%c                óf  — |j                  «       }|sd}n|r|dk(  sJ ‚d|›d�}nd|›�}| j                  rç| j                  d   j                  rÎ|j	                  «       r¾dt        | j                  «      › �}| j                  j                  t        ||› d|j                  |d¬	«      › �«      «       || j                  |<   t        j                  D ]E  }|j                  |«      }t        d
„ |D «       «      rŒ'| j                   |   }	||	vsJ d«       ‚||	|<   ŒG ||fS |j                  |«      }||fS )Nr‡   ú, other=0.0ú, boundary_check=z, padding_option='zero'r  r&  ú = F)r9  c              3  ó–   K  — | ]A  }t         j                  j                  j                  |t	        j
                  d «      «      –— ŒC y­w©r   N)rC   r  rù   rö   r®   r>  )r±   r—   s     rz   rÍ   z1TritonKernel.codegen_block_ptr.<locals>.<genexpr>  s9   è ø€ ò àô —G‘G×$Ñ$×<Ñ<¸VÄUÇ]Á]ÐSTÓEU×Vñùs   ‚AA	z@duplicate advancement for pointer '{block_ptr}' at type '{symt}')rL  r$  r,  rÀ  rÅ   Únextr   r›  r�   rL   rD  r¢  r–   r©   rP  r  r£  )
r¿   rA  rý  r)  rœ  Úcheckr&  r    Úadvance_offsetsÚadvancementss
             rz   Úcodegen_block_ptrzTritonKernel.codegen_block_ptrý  sp  € ð ×'Ñ'Ó)ˆÙà‰EÙØ˜MÒ)Ð)Ð)Ø'¨ yÐ0GÐH‰Eà'¨ yÐ1ˆEà×!Ò!Ø× Ñ  Ñ$×,Ò,Ø×#Ñ#Ô%à#¤D¨×):Ñ):Ó$;Ð#<Ð=ˆIØ�I‰I×ÑÜØ˜Y˜K s¨8¯?©?¸3È¨?Ó+NÐ*OÐPóôð 37ˆD×$Ñ$ YÑ/ô &×5Ñ5ò :�Ø"*×":Ñ":¸4Ó"@�ô ñ à"1ôô ð à#×8Ñ8¸Ñ>�Ø ¨Ñ4ð ØVóÐ4ð +:�˜YÒ'ð:ð" ˜%ÐÐð !Ÿ™¨Ó,ˆIØ˜%ÐÐr|   c                óÄ   — |j                  ||j                  |j                  d«      }|› dt        t        j
                  j                  |«      «      › d�}d|› d|› |› d�S )NFr  rü   ú	tl.store(rû   )r  rè   rí   rñ  rC   r  Ú	get_dtype)r¿   rA  r)  r&  r  rœ  s         rz   Úcodegen_block_ptr_store_linez)TritonKernel.codegen_block_ptr_store_line,  sj   € à×6Ñ6Ø�8×'Ñ'¨×)=Ñ)=¸uó
ˆð
 �'˜Ô/´·±×0AÑ0AÀ$Ó0GÓHÐIÈÐKˆØ˜9˜+ R¨ w¨u¨g°QÐ7Ð7r|   c                óÌ  — |s|sy t        |t        j                  «      sJ ‚| j                  |d¬«      }t        |t        «      sJ ‚|j
                  }|j                  «       r|j                  nd }|rt        | j                  |«      «      nd }| j                  ||rdnd ||«      }	| j                  |«      }
| j                  j                  |
|	dt        j                  ¬«       y )NFr%  Ú0)Ú
assignmentrè  )rb  r®   ÚExprr)  r³   rµ   rÀ   rÚ   ÚtexprÚrename_indexingÚindirect_assertÚget_load_bufferr-  r.  rs   rõ  )r¿   r  rj  Úlowerr°   r)  rµ   rÚ   Úsize_strÚlineÚbuffers              rz   Úcheck_boundszTritonKernel.check_bounds6  sÎ   € ñ ™Øä˜$¤§
¡
Ô+Ð+Ð+Ø—=‘= °�=Ó7ˆÜ˜(¤OÔ4Ð4Ð4à×&Ñ&ˆ	Ø(0×(9Ñ(9Ô(;�8×$Ò$ÀˆÙ8=”5˜×-Ñ-¨dÓ3Ô4À4ˆð ×#Ñ#Ø™e‘s¨¨x¸ó
ˆð ×%Ñ% hÓ/ˆØ�‰×Ñ˜& $°5ÄÇÁÐÕLr|   c                óô   — |j                  «       s|j                  «       r| j                  S | j                  r5| j                  d   j
                  r|j                  «       s| j                  S | j                  S )Nr  )	rÃ   rÑ   r/  r$  r,  rÀ  rÅ   r›  Úloads)r¿   r)  s     rz   r5  zTritonKernel.get_load_bufferP  sb   € Ø× Ñ Ô" h×&:Ñ&:Ô&<à—<‘<Ðà×!Ò!Ø× Ñ  Ñ$×,Ò,Ø×'Ñ'Ô)ð —9‘9Ðà—:‘:Ðr|   c           
     ó²	  ‡‡‡‡‡— | j                   j                  ‰«      }| j                  Š‰‰xx   dz  cc<   t        }| j	                  |«      Š|}| j                  |d¬«      }|j                  «       Š|j                  «       }t        d„ | j                  |«      j                  «       D «       «      }| j                  |«      rd}	nX|sd}	nS| j                  rE| j                  d   j                  r,ˆˆˆˆˆfd„}
‰‰   Šd}	t        j                   t"        d	|
«      }nd
}	|s‰r8|j%                  «       r(| j&                  rdt)        | j&                  «      › �}nd}nd
}	 d}t*        j,                  j.                  r"| j0                  j3                  «       }|‰   dkD  }	 | j                  |«       xr | j                   xr | xr |}d
}|rd}d }t4        j6                  j9                  ‰«      }t;        ‰«      r7|}|t<        j>                  t<        j@                  fv �rWt<        jB                  }�nEtE        |tF        «      rL| jI                  ‰|||«      \  }}d|› |› |	› |› d�}|jK                  ||jL                  |jN                  d«      }nVtE        |tP        jR                  «      rd|› d|› d�}|jT                  }n&d|› d|jV                  › d|jX                  › |	› |› |› d�
}|t<        j>                  t<        j@                  fv r/t*        j,                  jZ                  r|dz  }t<        jB                  }|t<        j\                  k(  r/t<        j^                  j`                  €|dz  }t<        j\                  }| jc                  |«      }| jd                  jg                  | ||«      |¬«      }|jh                  dkD  r‰‰xx   dz  cc<   tE        |tj        «      sJ ‚|jl                  |_6        |r±d|› d|› d�}| jd                  jg                  |||¬«      }|jl                  r~|jn                  rd}n|t<        j\                  k(  rd}nd}| j&                  rt)        | j&                  «      n|}d|jX                  › d|› d|› d�}| jd                  jg                  |||¬«      }| j                  r|jq                  «       s‰s| jr                  ju                  |«       |S )NrE   Tr%  c              3  ó&   K  — | ]	  }|d k(  –— Œ y­wrõ   r…   )r±   rŠ  s     rz   rÍ   z$TritonKernel.load.<locals>.<genexpr>t  s   è ø€ ò 
ØˆA��Fñ
ùó   ‚z, eviction_policy='evict_last'r  c                 ó    •— ‰‰   ‰ kD  r‰s‰ryy)NÚ
evict_lastÚevict_firstr…   )Úexpected_countrÅ   Úindirect_indexingÚload_countsrA  s   €€€€€rz   Údecide_laterz'TritonKernel.load.<locals>.decide_later}  s   ø€ Ø˜tÑ$ ~Ò5ÙÑ"3à'Ø$r|   z, eviction_policy='<EP>'z<EP>r‡   z, other=r   z, cache_modifier='.cg'rF  rü   r@  r´  r  r  r6  rç  rú   rû   z0.0r>  r/  r¦  );rC  rI  r¥  r   r  r)  rÅ   rÑ   rÐ   Úget_strides_of_loadr!  Úis_broadcastedr$  r,  rÀ  rä  rå  r3   rÀ   Ú_load_otherrV   r   rŒ   Úskip_l1_cacher&  Úbuffer_read_countsrC   r  r,  r_   rs   r  r  rt   rb  rß   r)  r  rí   rè   r®   r>  r¹   rµ   rÚ   r  r}   rx  ry  r5  r-  r.  Ú	use_countr   r·   rú  rÕ   rž  ra  )r¿   rA  r¼   rý  Ú	make_lineÚoriginal_indexr)  rÑ   Úis_coalescedÚeprF  rœ  Úhas_read_depsrK  rJ  ÚcachemodÚappend_broadcastrè  r8  r&  Úload_bufferÚ
result_varÚzeroÚ	other_valrC  rÅ   rD  rE  s    `                      @@@@rz   r9  zTritonKernel.load_  s¶  ü€ Ø�i‰i�o‰o˜dÓ#ˆØ×'Ñ'ˆØ�DÓ˜QÑÓÜCKˆ	Ø ×5Ñ5°eÓ<ÐØˆØ—=‘= °$�=Ó7ˆØ×(Ñ(Ó*ˆ
Ø×*Ñ*Ó,ˆô ñ 
Ø ×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ÜŸ™’ô ˜(¤OÔ4Ø#'×#9Ñ#9¸$ÀÀXÈuÓ#UÑ �	˜5Ø! ) ¨U¨G°B°4¸°zÀÐC�Ø×=Ñ=Ø˜(×.Ñ.°×0DÑ0DÀdó‘ô ˜N¬E¯M©MÔ:Ø! #  d¨>Ð*:¸"Ð=�Ø#+×#6Ñ#6Ñ à! #  d¨8×+=Ñ+=Ð*>¸cÀ(×BSÑBSÐATÐUWÐTXÐY^ÐX_Ð`hÐ_iÐijÐk�ð œ%Ÿ-™-¬¯©Ð8Ñ8Ü—M‘M×8Ò8àÐ)Ñ)�ÜŸ™�ØœŸ
™
Ò"¤u§}¡}×'8Ñ'8Ð'@ð ˜Ñ&�ÜŸ
™
�à×*Ñ*¨8Ó4ˆØ—X‘X×&Ñ& {±I¸d³OÈ5Ð&ÓQˆ
Ø×Ñ !Ò#Ø˜Ó Ñ"ÓÜ˜*Ô&7Ô8Ð8Ð8Ø'×1Ñ1ˆ
ÔáØ% j \°Ð4DÐ3EÀQÐGˆDØŸ™×*Ñ*¨;¸ÀEÐ*ÓJˆ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×.Ñ.¨{¸DÈÐ.ÓN�
à×$Ò$¨X×-?Ñ-?Ô-AÉ*Ø×"Ñ"×&Ñ& zÔ2àÐr|   c           	     ó¦  — | j                   j                  |«      }|}| j                  |d|d u ¬«      }|| j                   j                  v }| j	                  |«      }	|r'|	r%| j
                  j                  t        |d«      «       t        |t        «      r,| j                  |||«      \  }
}| j                  |||
||«      }n]|€$d|› d|j                  › d|› d|j                  › d�	}n7|d	k(  r$d
|› d|j                  › d|› d|j                  › d�	}nt        d|› �«      ‚t        j                   «       }| j"                  s7| j$                  r+|j'                  | j)                  || j
                  «      «       | j
                  j                  t        ||«      «       | j"                  s| j*                  j-                  |«       |j/                  «        y )NT)rÔ  r&  ztl.debug_barrier()r+  r@  r  rû   rü   Ú
atomic_addztl.atomic_add(z, sem='relaxed')zstore mode=)rC  r7  r)  Úinplace_buffersrH  Ústoresr�   rL   rb  rß   r)  r-  rµ   rÚ   r˜  Ú
contextlibÚ	ExitStackr$  r©  Úenter_contextÚguard_cooperative_storerž  ra  Úclose)r¿   rA  r¼   r  Úmoderý  rN  r)  Ú
is_inplacerH  r&  rœ  r8  Ú
exit_stacks                 rz   ÚstorezTritonKernel.storeæ  s·  € ð �i‰i×Ñ˜tÓ$ˆØˆØ—=‘= °tÀtÈtÀ|�=ÓTˆð ˜TŸY™Y×6Ñ6Ð6ˆ
Ø×,Ñ,¨^Ó<ˆÙ™.Ø�K‰K×!Ñ!¤,¨tÐ5IÓ"JÔKä�h¤Ô0Ø#×5Ñ5°d¸CÀÓJÑˆI�uà×4Ñ4Ø�h 	¨5°%ó‰Dð ˆ\Ø˜s˜e 4¨×(:Ñ(:Ð';¸3¸u¸gÀRÈ×HYÑHYÐGZÐZ[Ð\‰DØ�\Ò!Ø# C 5¨¨X×-?Ñ-?Ð,@ÀÀEÀ7È"ÈX×M^ÑM^ÐL_Ð_oÐp‰Dä%¨°D°6Ð&:Ó;Ð;ä×)Ñ)Ó+ˆ
Ø×$Ò$¨×)CÒ)CØ×$Ñ$ T×%AÑ%AÀ$ÈÏÉÓ%TÔUà�‰×Ñœl¨4°Ó6Ô7à×$Ò$Ø×"Ñ"×&Ñ& uÔ-à×ÑÕr|   c                ó”   — | j                   j                  «       }|j                  t        |d|› d�«      «       |j	                  «       S )z«
        For cooperative reductions only one thread block should write out the result.
        We rotate which thread block does each write for better parallelism
        zif rsplit_id == (z % RSPLIT):)r¹  r†  r�   rL   Úindent)r¿   rA  r9  rJ  s       rz   r_  z$TritonKernel.guard_cooperative_store	  sC   € ð
 ×8Ñ8×NÑNÓPˆØ×Ñœ dÐ.?À¸uÀKÐ,PÓQÔRØ�}‰}‹Ðr|   c                ó‚  — | j                   j                  t        j                  «       | j                  j                  |d   «      }| j                  |d   «      }	| j                  |d   «      }
| j                  |d   «      }|r| j                  j                  |d   «      nd}|r| j                  |d   «      nd}|t        j                  k(  rd}n!|t        j                  k(  rd}nt        d«      ‚| j                  j                  | j                  d	|› d
|› d
|	› d
|
› d
|› d
|› d
|› d
|› d
|› d
|› d
|› d�|¬«      }|S )z3
        See [Note: Inductor bucketize op]
        r   rE   r   r   r   r'  ztl.int64z5Bucketize only supports indexing with int32 and int64z'triton_helpers.bucketize_binary_search(rû   z, )rç  )r¦  ra  r(   ÚONE_ELEMENT_PER_THREADrC  rI  r  rs   rõ  r*  r˜  r-  r.  r/  )r¿   r!  Ú
boundariesÚboundary_indicesÚindexing_dtyper  ÚsorterÚsorter_indicesÚboundaries_ptrÚboundary_sizeÚboundaries_underlying_numelÚboundary_strideÚ
sorter_ptrÚsorter_striderA  r5  s                   rz   Ú	bucketizez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Ø5°f°X¸RØÐ˜b  ¨rÐ2MÐ1NÈbÐQ`ÐPaÐacØÐ  Øˆn˜BØˆg�RØˆl˜"˜]˜O¨2ØÐð ðð !ð #ó 
ˆð ˆr|   c                óž   — | j                  «       }|dk(  rd|› d�S | j                  }dg||z
  z  dg|z  z   }|› ddj                  |«      › d�S )	NrE   z!triton_helpers.promote_to_tensor(rü   r_  r   r`  rû   ra  )r"  r#  rØ   )r¿   r  ÚndimsÚnreduceÚsizess        rz   Úreduction_resizezTritonKernel.reduction_resizeL	  sj   € Ø×'Ñ'Ó)ˆØ�AŠ:Ø6°u°g¸QÐ?Ð?à×)Ñ)ˆØ�˜ ™Ñ)¨V¨H°wÑ,>Ñ>ˆØ�˜˜$Ÿ)™) EÓ*Ð+¨1Ð-Ð-r|   c           
     óô   — | j                   dk(  r|S | j                  «       | j                   z
  }| j                  «       }|d| dgz   }t        | j                  j                  |t        |||«      |¬«      «      S )zC
        Reshape to RBLOCK, collapsing all reduction dims.
        rE   NÚRBLOCKrç  )r#  r"  Údense_size_listr´   r-  r.  r   )r¿   r9  r  rè  Útarget_ndimr  Útarget_shapes          rz   Úreduction_collapse_dimsz$TritonKernel.reduction_collapse_dimsU	  s†   € ð
 ×"Ñ" aÒ'ØˆLà×-Ñ-Ó/°$×2IÑ2IÑIˆØ×,Ñ,Ó.ˆØ$ \ kÐ2°h°ZÑ?ˆÜØ�H‰H×ÑØœ u¨m¸\ÓJÐRWð ó ó
ð 	
r|   c                óÆ  ‡ ‡‡‡2‡3‡4‡5‡6‡7— d8d„}t        j                  |«      D �cg c]  }|j                  ‘Œ }}t        j                  ||«      }t	        d„ |D «       «      rHt        j                  |t
        j                  «      }t        j                  ‰t
        j                  «      Š‰ j                  sJ ‚t        d„ ‰ j                  D «       «      }‰ j                  |«       t        |«      }‰ j                  r|j                  ‰ j                  «       ‰ j                  d   j                  d   }	‰ j!                  «       Š3‰ j#                  ˆ3ˆ fd„|«      }‰ j%                  «       ‰ j&                  z
  Š4	 	 	 	 	 	 d9ˆ4ˆˆˆ fd„Š5	 	 	 	 	 	 	 	 d:ˆ5fd„}
ˆ4ˆˆ6ˆ fd	„}|‰|f}|‰ j(                  j*                  v r‰ j(                  j*                  |   S t-        |«      }t/        |«      }‰ j(                  j1                  |¬
«      }t        d„ |D «       «      |_        dj5                  |«      Š2ˆ2fd„Š7‰ j6                  �r!t8        j:                  j=                  ‰|«      }‰ j#                  t>        |«      }d;ˆ ˆ7fd„}‰dk(  rn?tA        |tB        «      r&tE        ||«      D ��cg c]  \  }} |||«      ‘Œ }}}n	 |||«      }‰dv r’tG        ‰ j(                  jI                  ‰ jJ                  d|	› d› d�tL        jN                  jP                  dk(  rt
        jR                  nt
        jT                  ¬
«      «      }dddœ‰   Š6 |‰ jJ                  |||«       �n™‰dk(  r8‰ jV                  r‰ jY                  |‰|‰7|‰«      }�np‰ j[                  ‰|«      }�n\‰dk(  rMtA        t\        «      sJ ‚|\  }}}tC        ˆˆ fd„‰ j_                  ‰ jJ                  |||‰4‰«      D «       «      }�n
‰dk(  r‰ ja                  ‰|«      }�nñtA        tb        «      sJ ‚‰ j(                  jI                  ‰ jJ                   ‰5‰ jJ                  tG        |«      d «      |j                  ¬
«      }�n�‰ j(                  je                  d|› �|¬
«      }t8        j:                  jg                  ‰|«      }‰ j#                  t>        |«      }tA        |tB        «      s5‰ jh                  jk                  |› d‰ j!                  «       › d|› d|› d�«       ‰dv rñd|› d�}‰ jl                  jo                  «       }‰ jh                  jk                  |› d‰ j!                  «       › dt        jp                  |«      jr                  › d‰ ju                  |«      › d�«       dddœ‰   Š6‰ jJ                  jw                  d|› d |› d!‰6› d"|› d|› d|› d|	› d#|› d$ ‰7|› d%�|«      › d&|› d$ ‰7|› d%�|«      › d&�«        |‰ jx                  |||«       �n t{        ‰«      r‰ jY                  |‰|‰7|‰«      }�nÝ‰dk(  �r(d|› d'�}d|› d(�}‰ jh                  jk                  |› d‰ j!                  «       › d)|› d�«       ‰ jh                  jk                  |› d*‰ j!                  «       › d|› d�«       ‰ jJ                  jw                  d+|› d |› d,|› d|› d|› dt|        j~                  › d-�«       ‰ jJ                  jw                  d+|› d$ ‰7|› d%�|«      › d+|› d$ ‰7|› d%�|«      › d+�	«       |}‰ j(                  j1                  ‰¬
«      }‰ j�                  ‰ jx                  ||||‰4‰«      }n¯t9        j‚                  ‰|«      } |||«      } ‰ jJ                  jk                  |› d$ ‰7| |«      › �«       |t
        j„                  k(  r/|› d.�}!t‡        ‰«      }" |
‰ jx                  tG        |«      |!|"«       n' |
‰ jx                  tG        |«      tG        |«      d «       ‰ jV                  �r{t8        j:                  jg                  ‰|«      }t‰        jŠ                  «       }#‰ jx                  ‰ jŒ                  fD ]2  }$|$jk                  d/«       |#j�                  |$j‘                  «       «       Œ4 ‰dv r©‰ jx                  jk                  |› d0‰ j“                  |› d1�«      › �«       ‰ j•                  |› d2�||«      }%‰ jl                  jo                  «       }‰ j•                  ||t        jp                  |«      jr                  «      }& |‰ jŒ                  ||%|&«       �n=t{        ‰«      rŽ‰dk(  sJ ‚|\  }'}(})‰ j•                  |'t/        |«      |d   «      }*‰ j•                  |(t/        |«      |d3   «      }+‰ j•                  |)t/        |«      |d4   «      },‰ j—                  ‰ jŒ                  |'|(|)|*|+|,‰4‰«	       n¤‰dk(  re|\  }}‰ j•                  |t/        |«      |d   «      }-‰ j•                  |t/        |«      |d3   «      }.‰ j�                  ‰ jŒ                  |||-|.‰4‰«       n:‰ j•                  |t/        |«      |«      }/ |
‰ jŒ                  tG        |«      |/d «       |#j™                  «        |‰ j(                  j*                  |<   tA        |tB        «      rÁt›        d5„ |D «       «      sJ ‚‰ jœ                  jŸ                  |«       ‰d6v rt¡        |«      d3k(  sJ ‚t¡        |«      |z  }t¡        |«      t¡        |«      k(  sJ ‚tE        ||«      D ]F  \  }0}1|1€J ‚|0j                  |1k7  sŒ‰ jx                  jk                  |0› d$|0› d7t‡        |1«      › d�«       ŒH |S tA        |t¢        «      sJ ‚‰ jœ                  j¥                  |«       |j                  |d   k7  r7|d   €J ‚‰ jx                  jk                  |› d$|› d7t‡        |d   «      › d�«       |S c c}w c c}}w )<Nc                ó¦   — | j                   t        j                  t        j                  fv r$t	        j
                  | t        j                  «      S | S rq   )rè  rs   r  r  rx  r<  rt   r:  s    rz   Úmaybe_upcastz,TritonKernel.reduction.<locals>.maybe_upcastm	  sF   € ð —;‘;ä—M‘MÜ—N‘Nðñô —‘˜U¤E§M¡MÓ2ðð ðr|   c              3  ó`   K  — | ]&  }|t         j                  t         j                  fv –— Œ( y ­wrq   )rs   r  r  r¯  s     rz   rÍ   z)TritonKernel.reduction.<locals>.<genexpr>}	  s"   è ø€ ÒM¸ˆq”U—]‘]¤E§N¡NÐ3Ô3ÑMùs   ‚,.c              3  ó:   K  — | ]  }|j                   › d �–— Œ y­w©rÌ   N©r!  rÁ  s     rz   rÍ   z)TritonKernel.reduction.<locals>.<genexpr>ƒ	  ó   è ø€ ÒM°D˜dŸk™k˜]¨$Ô/ÑMùó   ‚r  r   c                óv   •— ‰j                   j                  ‰j                  d| › d‰› d�| j                  ¬«      S )Nrú   rû   rü   rç  ©r-  r.  r/  rè  )r   r  r¿   s    €€rz   ú<lambda>z(TritonKernel.reduction.<locals>.<lambda>’	  s<   ø€ �d—h‘h×'Ñ'Ø—‘Ø" 1 # R¨Ð'7°qÐ9Ø—g‘gð (ó € r|   c           
     óÔ   •— ‰dv }|rdnd}‰j                  | |‰«      }‰dv r‰j                  |› d‰› d|› d‰› d�«      }n‰j                  |› d‰› d	|› d‰› d�«      }|�|› d
|› d�}|S )zK
            Helper to generate a reduction call, e.g. tl.sum.
            )rÐ   r7  ÚminÚprodÚtriton_helpersÚtl)r7  r�  rÝ  z2(rû   rü   r®  r  )r  ry  )	r9  r  Úresult_typeÚ
use_helperÚmoduler  rè  Úreduction_typer¿   s	        €€€€rz   Úfinal_reductionz/TritonKernel.reduction.<locals>.final_reduction�	  s±   ø€ ð (Ð+HÐHˆJÙ)3Ñ%¸ˆFà×0Ñ0°¸ÀÓFˆEØ Ñ/Ø×-Ñ-Ø�h˜a Ð/¨r°%°¸¸3¸%¸qÐAó‘ð ×-Ñ-Ø�h˜a Ð/¨q°°°r¸#¸¸aÐ@ó�ð Ð&Ø ˜'  k ]°!Ð4�àˆLr|   c                óF   •—  ‰| ||«      }| j                  |› d|› �«       y)zU
            Generate a reduction and assign it to an existing variable.
            r"  N)r�   )r9  rU  r  r‘  r•  s       €rz   Úfinal_reduction_definez6TritonKernel.reduction.<locals>.final_reduction_define·	  s*   ø€ ñ $ F¨E°;Ó?ˆEØ�M‰M˜Z˜L¨¨E¨7Ð3Õ4r|   c                óÊ   •— ‰j                  | |‰«      }‰j                  | |‰«      }| j                  d|› d|› d‰› d|› d|› d‰› d|› d‰j                  |› d�«      › d	�«       y )
Nú                z_val, z_idx = triton_helpers.z_with_index(rû   ú)
                r"  Ú_idxú
                )r  r�   ry  )r9  rU  r  r¼   r  rè  Úroot_opr¿   s       €€€€rz   Úfinal_argreducez/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ç  c              3  ó>   K  — | ]  }t        |d    «      rŒ|–— Œ y­wr$  )r9   )r±   rý  s     rz   rÍ   z)TritonKernel.reduction.<locals>.<genexpr>Ô	  s!   è ø€ ò *
ØÔ(;¸CÀ¹FÕ(CŒCñ*
ùs   ‚–r×   c                ó:   •— ‰s| S t         j                  ‰| |«      S rq   )r   rƒ  )ÚtvalÚfvalÚconds     €rz   Ú
where_condz*TritonKernel.reduction.<locals>.where_condÙ	  s    ø€ ÙØ�Ü(×.Ñ.¨t°T¸4Ó@Ð@r|   c                óv   •— ‰j                   j                  ‰j                   ‰| |«      | j                  ¬«      S )Nrç  rŠ  )r  Údefaultr¿   r¤  s     €€rz   Ú_mask_valuez+TritonKernel.reduction.<locals>._mask_valueâ	  s5   ø€ Ø—x‘x×(Ñ(Ø—L‘L¡*¨U°GÓ"<ÀEÇKÁKð )ó ð r|   Úonline_softmax_reduce)ÚargmaxÚargminrú   zindex, r	  r'  r7  r�  Úwelford_reduceÚwelford_combinec              3  ón   •K  — | ],  }‰j                   j                  ‰j                  |‰¬ «      –— Œ. y­w)rç  N)r-  r.  r/  )r±   r  rè  r¿   s     €€rz   rÍ   z)TritonKernel.reduction.<locals>.<genexpr>
  s3   øè ø€ ò #àð —H‘H×%Ñ% d§l¡l°EÀÐ%×Gñ#ùs   ƒ25rŒ  ú = tl.full(rû   rü   Ú_indexr™  ú_next, z_next = triton_helpers.z%imum_with_index(
                    z(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Ú_bvalrE   r   c              3  ó<   K  — | ]  }t        |t        «      –— Œ y ­wrq   )rb  r   r¯  s     rz   rÍ   z)TritonKernel.reduction.<locals>.<genexpr>Ð
  s   è ø€ ÒL¸A”z !Ô%6×7ÑLùrú  )r«  r¨  r  )r  rK   r~   rK   )r  r´   r‘  r¸   r~   r´   )rU  r´   r  r´   r‘  r¸   r~   r   )r~   rK   )SÚpytreeÚtree_leavesrè  Útree_maprÐ   rs   r1  rt   r$  r   r,  rÿ  Úsortedr‘  rc  r!  r  Ú_map_tuple_or_scalarr"  r#  r-  Úreduction_cacherø  rö  rM  r·   rØ   rÏ  r   Ú	ReductionÚdefault_valuerV   rb  r¶  rý   r´   r.  r/  rC   r  rr  rõ  r*  r©  r«  Úwelford_reduce_fallbackr   Ú_welfordÚ prepare_softmax_twopass_fallbackrK   ÚnamedvarÚdefault_accumulatorr›  r�   r&  Úselect_index_dtypeÚiinfor7  r¯  r�   rœ  r7   r   rc  Ú%online_softmax_reduce_final_reductionÚget_reduction_combine_fnr}   ré  r\  r]  r�  r^  rf  ry  Ú*codegen_cooperative_reduction_peer_combineÚwelford_reduce_final_reductionr`  r  rž  r  r  r   ra  )8r¿   rè  r2  r”  r  r‚  r3  Úoriginal_dtypesÚmasksÚreduction_range_prefixr—  rž  rO  Úacc_typeÚtorch_acc_typerU  r¦  r§  r   ÚdÚmasked_valueÚaccumulator_indexÚmeanÚm2ÚweightÚaccumulatorrr  Úaccumulator_maxÚaccumulator_sumÚ
result_maxÚ
result_sumÚ
combine_fnÚupdatedÚaccumulator_casted_strr‘  rc  Ú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  r  r•  r�  r¤  s8   `` `                                              @@@@@@rz   Ú	reductionzTritonKernel.reductionf	  s©  ÿø€ ó	ô 17×0BÑ0BÀ5Ó0IÖJ¨˜3Ÿ9›9ÐJˆÐJÜ—‘ ¨eÓ4ˆÜÑM¸_ÔMÔMä×+Ñ+¨I´u·}±}ÓEˆIÜ×'Ñ'¨¬u¯}©}Ó=ˆEà×$Ò$Ð$Ð$ÜÑM¸D×<LÑ<LÔMÓMˆØ×Ñ˜%Ô Ü�u“ˆØ�?Š?Ø�L‰L˜Ÿ™Ô)Ø!%×!1Ñ!1°"Ñ!5×!<Ñ!<¸QÑ!?Ðð ×,Ñ,Ó.ˆØ×)Ñ)ôð
 ó
ˆð ×%Ñ%Ó'¨$×*AÑ*AÑAˆð	àð	ð 'ð	ð ÷		ð 	ð4
	5àð
	5ð ð
	5ð 'ð	
	5ð
 õ
	5÷	ð  °Ð6ˆ	Ø˜Ÿ™×0Ñ0Ñ0Ø—8‘8×+Ñ+¨IÑ6Ð6ä" 9Ó-ˆÜ)¨)Ó4ˆØŸ(™(Ÿ/™/°˜/Ó?ˆ
Ü)ñ *
Ø ô*
ó  
ˆ
Ôð �z‰z˜%Ó ˆô	Að
 ×$Ó$Ü—l‘l×0Ñ0°ÀÓKˆGØ×/Ñ/´¸wÓGˆGöð Ð!8Ò8ð Ü˜E¤5Ô)Ü>AÀ%ÈÓ>Q×R±d°a¸¡¨A¨qÕ 1ÐR�ÒRá*¨5°'Ó:�àÐ!5Ñ5Ü$'Ø—H‘H×%Ñ%ØŸ™Ø*Ð+AÐ*BÀ'È,ÈÐW^Ð_äŸ8™8×/Ñ/°:Ò=ô $Ÿkškä"Ÿ[™[ð &ó ó%Ð!ð &+°eÑ<¸^ÑL�ÙØ—L‘L *¨lÐ<Möð  Ð#3Ò3Ø×-Ò-à!%×!4Ñ!4Ø" N°E¸:ÀxÐQVó"’Jð "&×!=Ñ!=¸eÀUÓ!K’JØÐ#4Ò4Ü! ,´Ô9Ð9Ð9Ø%1Ñ"��r˜6Ü"ô #à!%§¡ØŸ™ d¨B°¸¸Uó"ô#ó ’
ð  Ð#:Ò:ð "×BÑBÀ5È%ÓP’
ä! ,´Ô<Ð<Ð<Ø!ŸX™X×.Ñ.Ø—L‘LÙ# D§L¡L´#°lÓ2CÀTÓJØ&×,Ñ,ð /ó ’
ð Ÿ(™(×+Ñ+¨a°
¨|Ð,<ÀNÐ+ÓSˆKÜ—l‘l×6Ñ6°~ÀyÓQˆGØ×/Ñ/´¸wÓGˆGÜ˜g¤uÔ-Ø—	‘	×#Ñ#Ø"�m ;¨t×/BÑ/BÓ/DÐ.EÀRÈÀyÐPRÐS[ÐR\Ð\]Ð^ôð Ð!5Ñ5Ø&'¨
 |°6Ð$:Ð!Ø"Ÿm™m×>Ñ>Ó@�Ø—	‘	×#Ñ#Ø(Ð)¨°T×5HÑ5HÓ5JÐ4KÈ2Ü—{‘{ ;Ó/×3Ñ3Ð4°B°t×7HÑ7HÈÓ7UÐ6VÐVWðYôð &+°eÑ<¸^ÑL�à—‘×#Ñ#ðØ�˜WÐ%6Ð$7Ð7NÈwÈið XØ �M Ð$5Ð#6°b¸¸¸rÐBXÐAYð Zà�˜S¡¨{¨m¸5Ð,AÀ;Ó!OÐ Pð QØ"Ð# 3¡zÐ5FÐ4GÀuÐ2MÐO`Ó'aÐ&bð cðôñ  Ø×*Ñ*¨J¸ÐEVöô & nÔ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˜_Ó9�
à!×GÑGØ×*Ñ*ØØØ#Ø#ØØó‘
ô  ×8Ñ8¸ÈÓS�
Ù$ [°%Ó8�Ø—‘×&Ñ&Ø"�m 3¡z°'¸;Ó'GÐ&HÐIôð ¤§
¡
Ò*ð 1<¨}¸LÐ-IÐ*Ü"5°eÓ"<�KÙ*Ø×.Ñ.Ü˜J›Ø.Ø#õ	ñ +Ø×.Ñ.´°J³ÄÀ[ÓAQÐSWôð ×%Ó%Ü—l‘l×6Ñ6°~ÀyÓQˆGÜ#×-Ñ-Ó/ˆJØ×.Ñ.°×0DÑ0DÐEò 7�à—‘Ð.Ô/Ø×(Ñ(¨¯©«Õ6ð7ð
 Ð!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Ø%Ð)9Ò9Ð9Ð9Ø8BÑ5�˜Y¨Ø ×KÑKØÔ!1°)Ó!<¸gÀa¹jó�	ð ×IÑIØÔ/°	Ó:¸GÀA¹Jó�ð #×MÑMØ!Ô#3°IÓ#>ÀÈÁ
ó�ð ×3Ñ3Ø×(Ñ(ØØØ!ØØØØØõ
ð  Ð#:Ò:Ø)3Ñ&�
˜JØ×JÑJØÔ 0°Ó ;¸WÀQ¹Zó�ð  ×JÑJØÔ 0°Ó ;¸WÀQ¹Zó�ð ×:Ñ:Ø×(Ñ(ØØØØØØõð ×GÑGØÔ 0°Ó ;¸Wó�ñ 'Ø×(Ñ(¬#¨j«/¸5À$ôð ×ÑÔà.8ˆ�‰× Ñ  Ñ+ä�j¤%Ô(ÜÑ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õðð" Ðô ˜jÔ*;Ô<Ð<Ð<Ø×"Ñ"×&Ñ& zÔ2ð ×Ñ ?°1Ñ#5Ò5Ø& qÑ)Ð5Ð5Ð5Ø×&Ñ&×0Ñ0Ø!�l # j \°Ô6IÈ/ÐZ[ÑJ\Ó6]Ð5^Ð^_Ð`ôð Ðùò_ Kùód  Ss   ¥oÊoc                óš  — | j                  |||«      }| j                  |||«      }t        d«      D �cg c]'  }t        | j                  j	                  |¬«      «      ‘Œ) c}\  }}|j                  d|› d|› d|› d|› d|› dt        j                  › d|› d| j                  |› «      › d|› d| j                  |› «      › d�«       ||fS c c}w )Nr   rç  ú
            rû   z9 = triton_helpers.online_softmax_reduce(
                ú)
            r"  )	r  r*  r´   r-  rM  r�   r   rc  ry  )	r¿   r9  r×  rØ  r  rè  rŒ  rÙ  rÚ  s	            rz   Ú_online_softmax_reducez#TritonKernel._online_softmax_reduceì
  sù   € ð ×6Ñ6°v¸ÐPUÓVˆØ×6Ñ6°v¸ÐPUÓVˆÜMRÐSTËXÖ!VÈ¤# d§h¡h§o¡o¸E oÓ&BÕ"CÒ!VÑˆ
�JØ�‰ðØˆ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           	     óD  ‡ ‡‡— ˆˆˆ fd„|||fD «       \  }}}d|› d|› d|› d|› d�	}t        d«      D �cg c]'  }t        ‰ j                  j                  ‰¬«      «      ‘Œ) }	}‰j	                  dj                  |	«      › d|› �«       t        ˆ fd„|	D «       «      }
|
S c c}w )	z;
        Helper to codegen triton_helpers.welford.
        c              3  óD   •K  — | ]  }‰j                  ‰|‰«      –— Œ y ­wrq   )r  )r±   r  r9  rè  r¿   s     €€€rz   rÍ   z(TritonKernel._welford.<locals>.<genexpr>  s(   øè ø€ ò 
àð ×(Ñ(¨°¸×>ñ
ùs   ƒ ztriton_helpers.welford(rû   rü   r   rç  r"  c              3  ó@   •K  — | ]  }‰j                  |«      –— Œ y ­wrq   )ry  )r±   r  r¿   s     €rz   rÍ   z(TritonKernel._welford.<locals>.<genexpr>	  s   øè ø€ ÒX¸u˜d×3Ñ3°E×:ÑXùs   ƒ)r*  r´   r-  rM  r�   rØ   r¶  )r¿   r9  rÓ  rÔ  rÕ  r  rè  ÚwelfordrŒ  Úwelford_resultsÚresult_valuess   ``    `    rz   rÁ  zTritonKernel._welfordý
  s¬   ú€ õ
à  FÐ+ô
Ñˆˆb�&ð ,¨D¨6°°B°4°r¸&¸ÀÀCÀ5ÈÐJˆÜFKÈAÃhÖOÀœ3˜tŸx™xŸ™°U˜Ó;Õ<ÐOˆÐOØ×Ñ˜DŸI™I oÓ6Ð7°s¸7¸)ÐDÔEäÓXÈÔXÓXˆØÐùò	 Ps   ´,Bc                óÎ  — | j                  «       | j                  z
  }|› d�}|› d�}	|› d�}
| j                  j                  |› d| j	                  «       › d|› d�«       | j                  j                  |	› d| j	                  «       › d|› d�«       | j                  j                  |
› d| j	                  «       › d|› d�«       |dk(  r>|\  }}}| j
                  j                  d|› d	|	› d	|
› d
|› d|	› d|
› d|› d|› d|› d�«       n8|dk(  sJ ‚| j
                  j                  d|› d	|	› d	|
› d|› d|› d|	› d|
› d�«       | j
                  j                  d|› d ||› d�|«      › d|	› d ||	› d�|	«      › d|
› d ||
› d�|
«      › d�«       |}| j                  j                  |¬«      }| j                  j                  |¬«      }| j                  | j                  |||||	|
||«	      S )z%Helper to codegen a welford reductionÚ_meanÚ_m2Ú_weightr´  rû   rü   r¬  r™  r°  z<_next = triton_helpers.welford_combine(
                    z,
                    z#
                )
                r«  z;_next = triton_helpers.welford_reduce(
                    z1, roffset == 0
                )
                z            r"  r±  rí  rç  )r"  r#  r›  r�   r  r/  r�   r-  rM  rÊ  rœ  )r¿   rU  r”  r  r¤  rÎ  rè  r  rÖ  Úaccumulator_m2Úaccumulator_weightrÓ  rÔ  rÕ  rá  râ  rã  s                    rz   r«  zTritonKernel.welford_reduce  sª  € ð ×%Ñ%Ó'¨$×*AÑ*AÑAˆà#˜ EÐ*ˆØ&˜< sÐ+ˆØ *˜|¨7Ð3ÐØ�	‰	×ÑØˆ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ðô	
ð !ˆØ—H‘H—O‘O¨%�OÓ0ˆ	ØŸ™Ÿ™¨e˜Ó4ˆØ×2Ñ2Ø×"Ñ"ØØØØØØØØó

ð 
	
r|   c
                ó–   — | j                  ||||||	«      }
|||g}t        ||
«      D ]  \  }}|j                  |› d|› �«       Œ |||fS )z0Helper to codegen call to triton_helpers.welfordr"  )rÁ  rý   r�   )r¿   r9  rá  râ  rã  rÓ  rÔ  rÕ  r  rè  r!  Úresult_exprsÚresult_exprr  s                 rz   rÊ  z+TritonKernel.welford_reduce_final_reductionG  si   € ð —‘˜v t¨R°¸¸eÓDˆØ# Y°Ð>ˆÜ"% l°FÓ";ò 	6ÑˆK˜Ø�M‰M˜[˜M¨¨U¨GÐ4Õ5ð	6ð ˜I }Ð4Ð4r|   c                ó�   — | j                  |||||«      }||g}	t        |	|«      D ]  \  }
}|j                  |
› d|› �«       Œ ||fS ©Nr"  )rï  rý   r�   )r¿   r9  rÙ  rÚ  rç  rè  r  rè  r!  rý  rþ  r  s               rz   rÇ  z2TritonKernel.online_softmax_reduce_final_reduction[  se   € ð ×,Ñ,¨V°X¸xÈÈeÓTˆØ" JÐ/ˆÜ"% l°FÓ";ò 	6ÑˆK˜Ø�M‰M˜[˜M¨¨U¨GÐ4Õ5ð	6ð ˜:Ð%Ð%r|   c                óD   — | j                   r| j                   d   S t        S )NÚRSPLIT)r™  r+   r¾   s    rz   Ú
max_rsplitzTritonKernel.max_rsplite  s"   € Ø×ÒØ×$Ñ$ XÑ.Ð.Ü Ð r|   c                ó¾  — | j                   d   }| j                  «       sdnd}||j                  z  | j                  «       z  }| j                  j                  |«      \  }}| j                  j                  d|› d|› d| j                  |«      › dt        |«      › d|› d	|› d
|› d�d¬«       | j                  j                  |› d|› dt        |«      › d�«       |› d�S )a	  
        Generate code to save a [XBLOCK, RSPLIT] temporary workspace, where each thread block writes a different
        column.  After the barrier, every thread block loads the completed value so that it can compute the final
        value independently.
        r  zxindex < xnumelNrœ  z_ws = (rG  z).to(tl.pointer_type(z))
                tl.store(z%_ws + (xindex * RSPLIT + rsplit_id), rû   rî  T©Ústripz_peers = tl.load(z_ws + (xindex * RSPLIT + rsplit_arange), rsplit_mask, eviction_policy='evict_first', other=triton_helpers.if_mask(rsplit_mask, r´  Ú_peers)r%  r¼  rë  r  r¹  r€  rœ  r�   r  r=   r�  r�   rV   )	r¿   rU  rè  Údefault_valÚxnumelrÌ   r|  r~  r  s	            rz   rÉ  z7TritonKernel.codegen_cooperative_reduction_peer_combinej  s$  € ð —‘˜SÑ!ˆØ(,×(@Ñ(@Ô(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Ð+¨J¨<ð 8eÜerÐs~Óeð  eAð  ACðDô	
ð �˜VÐ$Ð$r|   c                óÐ  — | j                   sJ ‚d| _         | j                  |d¬«      }d| _         | j                  j                  |«      }t	        j
                  «       }| j                  r+|j                  | j                  || j                  «      «       t        |t        «      rY| j                  j                  t        || j                  |||j                  |«      |d|j!                  «       ›�«      «      «       nXt        |t"        «      sJ ‚| j                  j                  t        |d|› d|j$                  › d|› d|j&                  › d	�	«      «       |j)                  «        y )
NFTr%  r!  r+  r@  r  rû   rü   )r$  r)  rC  r7  r\  r]  r©  r^  r_  r�  rb  rß   r�   rL   r-  rD  rL  r³   rµ   rÚ   r`  )r¿   rA  r¼   r  r)  rý  rc  s          rz   Ústore_reductionzTritonKernel.store_reduction…  sR  € ð ×$Ò$Ð$Ð$Ø %ˆÔØ—=‘= °$�=Ó7ˆØ $ˆÔØ�i‰i×Ñ˜tÓ$ˆä×)Ñ)Ó+ˆ
Ø×%Ò%Ø×$Ñ$Ø×,Ñ,¨T°4×3GÑ3GÓHôô �h¤Ô0Ø× Ñ ×*Ñ*ÜØØ×5Ñ5ØØ Ø Ÿ™¨Ó,ØØ+¨H×,CÑ,CÓ,EÐ+HÐIóó	õô ˜h¬Ô8Ð8Ð8Ø× Ñ ×*Ñ*ÜØØ ˜u D¨×);Ñ);Ð(<¸CÀ¸wÀbÈ×IZÑIZÐH[Ð[\Ð]óôð 	×ÑÕr|   c           	     ó   ‡‡‡
‡‡‡‡— t        «       Š‰j                  d«       t        «       Š
t        d«      D �‡cg c]!  Št	        ˆ
ˆˆfd„t        |«      D «       «      ‘Œ# }}dj                  d„ t        j                  j                  |«      D «       «      }‰j                  d|› d�«       t        «       ŠdŠd	d
l
m}  |«       Š G ˆ
ˆˆˆˆfd„dt        «      }‰j                  «       5  t        j                   |«       «      5   ||Ž }	dj                  d„ |	D «       «      }	‰j                  d|	› �«       d d d «       d d d «       | j                   j#                  ‰j%                  «       ‰¬«      S c c}w # 1 sw Y   ŒAxY w# 1 sw Y   ŒExY w)Nz@triton.jitr   c              3  óV   •K  — | ]   }‰j                  d ‰› d|› �‰|   ¬«      –— Œ" y­w)r	  rŒ  rç  N)rÃ  )r±   Únr-  ÚdtypesrŠ  s     €€€rz   rÍ   z,TritonKernel._lift_helper.<locals>.<genexpr>·  s.   øè ø€ ÒXÀ1�#—,‘,  Q C q¨¨˜}°F¸1±I�,×>ÑXùs   ƒ&)rû   c              3  ó2   K  — | ]  }t        |«      –— Œ y ­wrq   ©r´   r¯  s     rz   rÍ   z,TritonKernel._lift_helper.<locals>.<genexpr>º  s   è ø€ ÒR¨œc !ŸfÑRùó   ‚zdef {name}(ú):rZ  r   rd   c                  ó2   •— e Zd Z	 	 	 	 	 	 	 	 dˆ ˆˆˆˆfd„Zy)ú+TritonKernel._lift_helper.<locals>.CSEProxyc                ó€   •— ‰d|› �z  Š t        ‰|«      |i |¤Ž}‰j                  ‰ t        ‰	|«      |i |¤Ž|¬«      S )NrŒ  rç  )rí  r.  )
r¿   rA  rC  r  Úoutput_dtyper-  Údtype_handlerÚhelperÚhelper_nameÚ	overridess
        €€€€€rz   Ú_defaultz4TritonKernel._lift_helper.<locals>.CSEProxy._defaultÊ  sq   ø€ ð   4 &˜zÑ)�ð œwØ!Øó ð ð #ð "ñ #�ð
 —|‘|ØØ,”G˜I tÓ,¨dÐ=°fÑ=Ø&ð $ó ð r|   N)rA  r´   rC  ztuple[Any, ...]r  údict[str, Any]r~   r   )rr   r€   r�   r  )r-  r  r  r  r  s   €€€€€rz   ÚCSEProxyr  É  s-   ø„ ðØðØ'6ðØ@Nðà÷ò r|   r  c              3  ó2   K  — | ]  }t        |«      –— Œ y ­wrq   r  )r±   r7  s     rz   rÍ   z,TritonKernel._lift_helper.<locals>.<genexpr>Ý  s   è ø€ ÒB°¤ F§ÑBùr  úreturn r[  )rM   r�   rJ   r*  r¶  rØ   r  r   Úfrom_iterabler-  r  re   r%   rf  rC   Úset_ops_handlerr•  ra  r‘   )r¿   ÚfnÚnum_argsr  rŠ  rC  Ú	signaturere   r  Úoutputsr-  r  r  r  r  s      ``     @@@@@rz   Ú_lift_helperzTritonKernel._lift_helper¯  s]  þ€ ô  Ó!ˆØ×Ñ˜Ô'Ü‹eˆô ˜1“X÷
àô ÕXÌÈhËÔXÕXð
ˆð 
ð —I‘IÑR¬i¯o©o×.KÑ.KÈDÓ.QÔRÓRˆ	Ø×Ñ˜=¨¨°2Ð6Ô7ä#Ó%ˆ	ð *ˆåPá2Ó4ˆ÷	ò 	”~ô 	ð$ �]‰]‹_ñ 	2œa×/Ñ/±³
Ó;ñ 	2Ù˜$�iˆGØ—i‘iÑB¸'ÔBÓBˆGØ×Ñ˜w w iÐ0Ô1÷	2÷ 	2ð
 ×$Ñ$×(Ñ(¨¯©Ó):ÀkÐ(ÓRÐRùòU
÷J	2ð 	2ú÷ 	2ð 	2ús)   »&E3Ã+FÄ2E8Ä8FÅ8F	Å=FÆFc                óâ  ‡ ‡— ‰ j                   sJ ‚‰ j                  rJ d«       ‚t        d„ ‰ j                  D «       «      }‰ j	                  |«       t        |«      }‰ j                  rJ d«       ‚g }g }t        d„ |D «       «      }t        j                  ‰ j                  j                  ‰ j                  «      }‰ j                  |t        |«      |«      }‰ j                  «       ‰ j                   z
  }	t#        ||«      D �]1  \  }
}‰ j                  j                  ‰ j                  |
› dt%        |«      › d�|¬«      }‰ j                  j                  ‰ j                  d|› d	‰ j'                  «       › d�|¬«      }
|j)                  |
«       t+        |«      }‰ j,                  rŒ¤‰ j                  j/                  |¬«      }‰ j1                  «       }d
|d<   dd	j3                  |«      › d�}|j4                  rdnd}‰ j6                  j9                  |› d|› d	|› d	|› d�«       |j)                  |«       �Œ4 d„ Šˆˆ fd„} |d ‰|«      › d|	› d	|› d�|||«      }‰ j,                  sÍ|D �cg c]#  } |d|› d�t;        |j<                  «      ¬«      ‘Œ% }} |t        |«      t        |«      «      } |t        |«      |«      }t#        ||«      D ��cg c]   \  }} |d|› d	|› d�|j<                  ¬«      ‘Œ" }}}t#        |||«      D ]*  \  }}}‰ j                  j9                  |› d|› d	|› d�«       Œ, n|}|D ]$  }t?        |t@        «      sJ ‚t        |«      |_!        Œ& t        |«      S c c}w c c}}w )NÚTODOc              3  ó:   K  — | ]  }|j                   › d �–— Œ y­wr…  r†  rÁ  s     rz   rÍ   z$TritonKernel.scan.<locals>.<genexpr>ì  r‡  rˆ  z(ops.scan not supported inside ops.maskedc              3  ó2   K  — | ]  }t        |«      –— Œ y ­wrq   ©r?   ©r±   rè  s     rz   rÍ   z$TritonKernel.scan.<locals>.<genexpr>ô  ó   è ø€ ÒF°eÔ*¨5×1ÑFùr  r  rü   rç  rú   rû   r^  r  r`  ra  zfloat('nan')z-1r®  c                ó2   — dj                  d„ | D «       «      S )Nr³  c              3  ó&   K  — | ]	  }|› d �–— Œ y­w©ú,Nr…   ©r±   r  s     rz   rÍ   z1TritonKernel.scan.<locals>.csv.<locals>.<genexpr>  ó   è ø€ Ò<¨E˜u˜g QœKÑ<ùr?  ©rØ   ©r!  s    rz   ÚcsvzTritonKernel.scan.<locals>.csv  ó   € Ø—8‘8Ñ<°VÔ<Ó<Ð<r|   c                ó  •— t        |«      }t        |«      D �cg c]  }| › d|› d|› �‘Œ }}t        ˆfd„|D «       «      r)|D �cg c]  }‰j                  j	                  |«      ‘Œ c}S |D �cg c]  }‰j                  j                  |¬«      ‘Œ  }	}‰j                  j                   ‰|	«      › d| › �«       t        |	|«      D ]*  \  }
}|r||
_	        ‰j                  j                  ||
«       Œ, t        |	«      S c c}w c c}w c c}w )Nrû   c              3  óT   •K  — | ]  }‰j                   j                  |«      –— Œ! y ­wrq   ©r-  Úcontains©r±   rO  r¿   s     €rz   rÍ   z:TritonKernel.scan.<locals>.cse_multiple.<locals>.<genexpr>  ó    øè ø€ ÒL°I�4—8‘8×$Ñ$ Y×/ÑLùó   ƒ%(rç  r"  )r  r*  r  r-  r^  rM  r/  r�   rý   r·   rN  r¶  )r8  r!  rÌ  r  r  rŠ  Ú
cache_keysrO  Ú_dtypeÚresult_varsrU  r7  r¿   s              €€rz   Úcse_multiplez'TritonKernel.scan.<locals>.cse_multiple  sû   ø€ Ü�F“ˆAÜ;@À»8ÖD°a˜T˜F " Q C r¨%¨Ò1ÐDˆJÐDÜÓLÀÔLÔLØAKÖL°I˜Ÿ™Ÿ™ YÕ/ÒLÐLØGMÖN¸V˜4Ÿ8™8Ÿ?™?°˜?Õ8ÐNˆKÐNØ�L‰L×"Ñ"Ù�{Ó#Ð$ C¨ vÐ.ôô *-¨[¸*Ó)Eò 4Ñ%�
˜IÙØ+0�JÔ(Ø—‘—‘˜Y¨
Õ3ð4ô ˜Ó%Ð%ùò EùâLùÚNs   šC=Á"DÁ.#Dztl.associative_scan((r  ztriton_helpers.select_one((z1), rbase == (RBLOCK - 1), dim=-1, keep_dims=True)ztl.where(roffset > 0, z = tl.where(roffset > 0, )"r$  r©  r   r,  rÿ  r»  r‘  r¶  rä  rå  r-  r.  r/  r'  r  r"  r#  rý   ré  r  rc  rø  rÏ  rM  r|  rØ   rú  r›  r�   r?   rè  rb  r   r·   )r¿   r  rÛ  r!  rÌ  Úbroadcasted_valuesÚaccumulatorsÚcse_computeÚcombine_helper_fnr  r  rè  Úvalue_dtyperÎ  rÖ  Úreduced_sizer¦  rC  Úpartial_scan_varsÚpartial_scan_varÚpartial_reduce_varsÚ	accs_nextÚfull_scan_varsÚ	full_scanÚpartial_scanrB  Úacc_nextÚpartial_reducerU  r7  s   `                            @rz   ÚscanzTritonKernel.scanâ  sÆ  ù€ ð ×$Ò$Ð$Ð$Ø×-Ò-Ð5¨vÓ5Ð-ÜÑM¸D×<LÑ<LÔMÓMˆØ×Ñ˜%Ô Ü�u“ˆØ—?’?ÐNÐ$NÓNÐ"àÐØˆäÑF¸vÔFÓFˆÜ×'Ñ'¨¯©×(9Ñ(9¸4¿<¹<ÓHˆØ ×-Ñ-¨j¼#¸f»+ÀvÓNÐØ×%Ñ%Ó'¨$×*AÑ*AÑAˆä ¨Ó/ó 	1‰LˆE�5ØŸ(™(×+Ñ+Ø—‘Ø�'˜Ô1°%Ó8Ð9¸Ð;Øð ,ó ˆKð
 —H‘H×%Ñ%Ø—‘Ø" ; -¨r°$×2EÑ2EÓ2GÐ1HÈÐJØð &ó ˆEð
 ×%Ñ% eÔ,ä& uÓ-ˆHà×,Ó,Ø"Ÿh™hŸo™o°E˜oÓ:�Ø#×3Ñ3Ó5�Ø#&�˜RÑ Ø!" 4§9¡9¨\Ó#:Ð";¸1Ð=�à,1×,CÒ,C™.È�Ø—	‘	×#Ñ#Ø"�m ;¨|¨n¸B¸w¸iÀrÈ(ÈÐSTÐUôð ×#Ñ# KÖ0ð5	1ò8	=õ	&ñ )Ø#¡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÷ñ
 ,�I˜|ñ	 Ø,¨Y¨K°r¸,¸ÀqÐIØ&×,Ñ,öðˆKñ ô :=Ø˜<Ð)<ó:ò Ñ5�˜+ ~ð —‘×&Ñ&Ø"�mÐ#<¸X¸JÀbÈÐHXÐXYÐZõñð ,ˆKà%ò 	5ˆJÜ˜jÔ*;Ô<Ð<Ð<Ü#-¨eÓ#4ˆJÕ ð	5ô �[Ó!Ð!ùò;#ùós   É'(M&Ë%M+c                ó  ‡ ‡— ‰ j                   sJ ‚‰ j                  rJ d«       ‚t        d„ ‰ j                  D «       «      }‰ j	                  |«       t        |«      }‰ j                  rJ d«       ‚‰ j                  sJ d«       ‚t        j                  ‰ j                  j                  ‰ j                  «      }‰ j                  «       ‰ j                  z
  }t        d„ |D «       «      }t!        |«      t!        |«      k(  sJ ‚t#        |«      D ��	cg c]'  \  }}	 |d|	› d‰ j%                  «       › d�||   ¬	«      ‘Œ) }
}}	d
„ Šˆˆ fd„}‰ j                  d   j&                  sJ ‚‰ j)                  ‰ j                  d   «      rdnd}t!        |«      dk(  r0d|
d   › d|
d   › d|› d|› d|› d|› d�} ||t!        |«      ||«      }nt+        d«      ‚t-        ||«      D ]  \  }}||_        |j0                  |_        Œ t        |«      S c c}	}w )Nr)  c              3  ó:   K  — | ]  }|j                   › d �–— Œ y­wr…  r†  rÁ  s     rz   rÍ   z$TritonKernel.sort.<locals>.<genexpr>Z  r‡  rˆ  z(ops.sort not supported inside ops.maskedz3ops.sort is only supported in persistent reductionsc              3  ó2   K  — | ]  }t        |«      –— Œ y ­wrq   r,  r-  s     rz   rÍ   z$TritonKernel.sort.<locals>.<genexpr>e  r.  r  rú   rû   rü   rç  c                ó2   — dj                  d„ | D «       «      S )Nr³  c              3  ó&   K  — | ]	  }|› d �–— Œ y­wr1  r…   r3  s     rz   rÍ   z1TritonKernel.sort.<locals>.csv.<locals>.<genexpr>o  r4  r?  r5  r6  s    rz   r7  zTritonKernel.sort.<locals>.csvn  r8  r|   c                ó  •— t        |«      D �cg c]  }| › d|› d|› �‘Œ }}t        ˆ
fd„|D «       «      r)|D �cg c]  }‰
j                  j                  |«      ‘Œ c}S t        |«      D �cg c]!  }‰
j                  j	                  ||   ¬«      ‘Œ# }}‰
j
                  j                   ‰	|«      › d| › �«       t        ||«      D ]*  \  }}|r||_        ‰
j                  j                  ||«       Œ, t        |«      S c c}w c c}w c c}w )Nrû   c              3  óT   •K  — | ]  }‰j                   j                  |«      –— Œ! y ­wrq   r;  r=  s     €rz   rÍ   z:TritonKernel.sort.<locals>.cse_multiple.<locals>.<genexpr>s  r>  r?  rç  r"  )r*  r  r-  r^  rM  r/  r�   rý   r·   rN  r¶  )r8  r  rÌ  r  rŠ  r@  rO  rB  rU  r7  r¿   s            €€rz   rC  z'TritonKernel.sort.<locals>.cse_multipleq  sú   ø€ Ü;@À»8ÖD°a˜T˜F " Q C r¨%¨Ò1ÐDˆJÐDÜÓLÀÔLÔLØAKÖL°I˜Ÿ™Ÿ™ YÕ/ÒLÐLÜEJÈ1ÃXÖNÀ˜4Ÿ8™8Ÿ?™?°¸±˜?Õ;ÐNˆKÐNØ�L‰L×"Ñ"Ù�{Ó#Ð$ C¨ vÐ.ôô *-¨[¸*Ó)Eò 4Ñ%�
˜IÙØ+0�JÔ(Ø—‘—‘˜Y¨
Õ3ð4ô ˜Ó%Ð%ùò EùâLùÚNs   �C>º"DÁ,&Dr  r   Úrnumelr   ztriton_helpers.sort_with_index(r   rE   z	, stable=z, descending=zUnhandled sort)r$  r©  r   r,  rÿ  r»  r‘  rÏ  rä  rå  r-  r.  r/  r"  r#  r¶  r  Ú	enumerater  r'  rµ  rz  rý   r·   r  )r¿   r  r!  ÚstableÚ
descendingrÌ  rF  r  rŠ  r  rD  rC  r[  r8  rB  rU  Ú	input_varr7  s   `                @rz   ÚsortzTritonKernel.sortQ  s2  ù€ ð ×$Ò$Ð$Ð$Ø×-Ò-Ð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ñ Ø" 5 '¨¨D×,?Ñ,?Ó,AÐ+BÀ!ÐDÈFÐSTÉIöð
Ðñ 
ò	=õ	&ð ×Ñ Ñ#×0Ò0Ð0Ð0Ø×2Ñ2°4×3CÑ3CÀBÑ3GÔH‘Èhˆäˆv‹;˜!Òà1Ð2DÀQÑ2GÐ1HÈÐK]Ð^_ÑK`ÐJað bØ�8˜2˜c˜U )¨F¨8°=ÀÀÈAðOð ñ ' t¬S°«[¸%ÀÓH‰Kä Ð!1Ó2Ð2ä%(¨°fÓ%=ò 	1Ñ!ˆJ˜	Ø#(ˆJÔ Ø )× 0Ñ 0ˆJÕð	1ô �[Ó!Ð!ùóQ
s   Ä,Hc                óø  — | j                   s=| j                  s1| j                  s%| j                  s| j                  s| j
                  sy| j                  D �cg c]  }|j                  sŒ|‘Œ }}| j                  �rWt        |«      dkD  �rHt        |«      D ]Ó  \  }}| j                  j                  |¬«      5  |j                  }| j                  rdnd}| j                  rdn|› d�}| j                  j                  d|› d	|› d
|› d
|j!                  «       › d�	«       ddd«       | j                  j                  |dz   ¬«      5  | j#                  || j                  «       ddd«       ŒÕ | j                  j                  t        |«      ¬«      5  | j%                  | j                  «       | j                  j'                  | j                   «       | j                  j'                  | j                  «       | j                  j'                  | j                  «       | j                  j'                  | j                  «       ddd«       t)        g t        |«      ¢«      D �]o  \  }}| j                  j                  |dz   ¬«      5  | j*                  |j,                     j/                  «       D ]á  \  }}|t        |«      dz
  k  rs||dz      }	| j*                  |	j,                     |   }
t0        j3                  |	«      }t5        |	j6                  |«      }t9        ||
«      D ��cg c]  \  }}|||z  z
  ‘Œ }}}| j                  j                  t;        | j<                  |   |› d|› d
t>        j@                  jC                  |«      › d�«      «       Œã 	 ddd«       | jD                  jG                  | jH                  «       |jK                  «        �Œr n”| j                  j'                  | j                   «       | j                  j'                  | j                  «       | j                  j'                  | j                  «       | j                  j'                  | j                  «       | j                  j'                  | j                  «       | j                  rb| j                  s| j
                  rJ| jL                  › d�}| j                  j'                  d|› d�d¬«       | jN                  jQ                  «        | j                  j'                  | j
                  «       | j                   jS                  «        | j                  jS                  «        | j                  jS                  «        | j                  jS                  «        | j                  jS                  «        | j
                  jS                  «        yc c}w # 1 sw Y   �Œ¥xY w# 1 sw Y   �ŒBxY w# 1 sw Y   �ŒœxY wc c}}w # 1 sw Y   �Œ^xY w)a  
        Concat output code from index_code, loads, compute, stores,
        suffix into self.body.

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

        For reduction kernels, this generates a loop over the reduction
        axis.
        Nr   )r—   Úrsplit_startr/  Ú
rsplit_endrÛ  zfor zoffset in range(rû   zBLOCK):rE   z = tl.advance(rü   z + tl.program_id(1)zR
                if HAS_RSPLIT:
                    triton_helpers.x_grid_barrier(rš  Tr  )*Úindexing_coder<  r[  r/  rœ  r�  r,  rÀ  r$  r  r\  r›  rf  r!  r©  r�   r°   rÃ  rÇ  r�   r)  r£  r    r  r–   r£   r   rÛ  rý   rL   r¢  rC   r  r  r-  Ú
invalidaterž  Úcache_clearr¸  r¹  rƒ  Úclear)r¿   r¢   Ú
loop_treesÚlevelr!  Ú
loop_startÚloop_endr&  ÚadvancementÚ	prev_treeÚprev_advancementÚ
prev_blockÚprev_num_iterÚcurÚprevÚsem_ptrs                   rz   Úcodegen_bodyzTritonKernel.codegen_body‘  s	  € ð ×ÒØ�zŠzØ�{Š{Ø�|Š|Ø×%Ò%Ø×#Ò#àà'+×'7Ñ'7ÖH˜t¸4¿<»<’dÐHˆ
ÐHØ× Ó ¤S¨£_°qÓ%8ä(¨Ó4ò J‘��tØ—Y‘Y×%Ñ%¨UÐ%Ó3ñ Ø!Ÿ[™[�FØ37×3MÒ3M¡ÐSV�Jà(,×(BÒ(B™È6È(ÐRWÐHXð ð —I‘I×'Ñ'Ø˜v˜hÐ&6°z°lÀ"ÀXÀJÈbÐQW×Q]ÑQ]ÓQ_ÐP`Ð`gÐhô÷ð —Y‘Y×%Ñ%¨U°Q©YÐ%Ó7ñ JØ×8Ñ8¸¸t¿y¹yÔI÷Jð JðJð —‘×!Ñ!¬¨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 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÷+á$- C¨ð !$ d¨]Ñ&:Ó :ð+˜Kñ +ð
 Ÿ	™	×+Ñ+Ü(Ø $× 8Ñ 8¸Ñ CØ#, +¨^¸I¸;ÀbÌÏÉ×I^ÑI^Ð_jÓIkÐHlÐlmÐ nóõñ!÷ð4 —‘×#Ñ# D×$:Ñ$:Ô;Ø× Ñ Ö"ñ9#ð< �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÷ñ ú÷Jñ Jú÷.ñ .üó,+÷ñ úsR   ÁV=Á*V=Â9A&WÅWÆB0WÊBW/Ì#W)Ì5AW/×W	×W	×W&×)W/×/W9	c                óV  — g }| j                  «       �rg }| j                  d|g «       |D ]û  }t        |t        «      r|j	                  t        |«      «       Œ.t        |t        «      rL|j	                  t        t        j                  j                  j                  |j                  «      «      «       ŒŠt        |t        j                  «      rB|j	                  t        t        j                  j                  j                  |«      «      «       Œæt        dt        |«      › �«      ‚ |S )Nr‡   z!Unsupported numel argument type: )rË  Úadd_numel_to_call_argsrb  r6  rc  r´   ra   rC   r  rù   Ú	size_hintÚ
inner_exprr®   r1  rà  rµ  )r¿   rC  Ú
numel_argsr	  s       rz   Úkernel_benchmark_extra_argsz(TritonKernel.kernel_benchmark_extra_argsõ  sÛ   € ØˆØ×ÑÕ!Ø+-ˆJØ×'Ñ'¨¨J¸Ô;Ø!ò V�Ü˜c¤3Ô'Ø—K‘K¤ C£Õ)Ü ¤_Ô5Ø—K‘K¤¤A§G¡G×$4Ñ$4×$>Ñ$>¸s¿~¹~Ó$NÓ OÕPÜ ¤U§Z¡ZÔ0Ø—K‘K¤¤A§G¡G×$4Ñ$4×$>Ñ$>¸sÓ$CÓ DÕEä$Ð'HÌÈcËÈÐ%TÓUÐUðVð ˆr|   c                ó  — t        «       }| j                  j                  «       \  }}}}|j                  g d¢«       |j	                  «       5  t        j                  «       }g }t        ||«      D �]È  \  }	}
dt        |«      › �}t        j                  j                  |	«      }|r¨|j                  |› dt        j                  j                  j                  |j                  «       «      › dt        j                  j                  j                  |j!                  «       «      › d|j#                  «       › d|j%                  «       › d�
«       �nÚ|	t        j                  j&                  v r½t        j                  j&                  |	   }|j                  |› dt        j                  j                  j                  |j)                  «       «      › dt        j                  j                  j                  |j+                  «       «      › d|j,                  › d|j.                  › d�
«       �nt1        |
t2        «      rZt        j                  j                  j5                  |
j6                  «      }d|
j8                  v rd	}|j                  |› d
|› �«       n—t1        |
t:        «      ryt        j                  j=                  «       }t        j                  j                  j5                  |
j                  «      }|j                  |› d|› d|› d|
j.                  › d�«       nt?        d|	› �«      ‚|jA                  |«       �ŒË |jC                  | jE                  «       «       |j                  ddjG                  |«      › d�«       d d d «       |j                  g d¢«       t        j                  j=                  «       }|jH                  }|j	                  «       5  |j                  dt        j                  jJ                  jM                  |«      › d�«       |j	                  «       5  |j                  t        j                  jJ                  jO                  |«      «       d|› �}|j                  |› d|› d�«       |j                  tQ        tR        jT                  «      › d|› d�«       d d d «       d d d «       |j                  g d¢«       |j	                  «       5  |j                  dt        j                  jJ                  jM                  |«      › d�«       |j	                  «       5  |j                  t        j                  jJ                  jO                  |«      «       |j                  dtQ        tR        jT                  «      › d�«       d d d «       d d d «       |j                  g d¢«       |j	                  «       5  |j                  d«       |j                  d«       |j                  d«       |j                  d«       |j                  d|› �«       |j                  d«       |j                  d«       d d d «       |S # 1 sw Y   �ŒÄxY w# 1 sw Y   �Œ¬xY w# 1 sw Y   �Œ±xY w# 1 sw Y   ŒàxY w# 1 sw Y   ŒäxY w# 1 sw Y   |S xY w)N)r‡   r‡   zdef get_args():Úarg_z = rand_strided(rû   z
, device='z	', dtype=rü   rJ  r   r"  z = torch.zeros(z*Don't find the buffer or const tensor for r   r2  )ú
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()z:ms = benchmarker.benchmark_gpu(lambda: call(args), 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"))+rM   rC  Úpython_argdefsÚ
writelinesrf  r  rŸ  rý   r%  rC   r  Útry_get_bufferr�   rù   Ú
size_hintsÚget_sizeÚ
get_strideÚ
get_devicer,  Ú	constantsrj  r.  Údevicerè  rb  rR   rw  r  rA  rT   Úget_current_device_or_throwÚKeyErrorrc  Úextendrz  rØ   r¼   Ú
device_opsÚdevice_guardÚ
set_devicer´   r8   ÚKERNEL_NAME)r¿   Únum_gbr5  Ú_argdefsÚ	call_argsr%  rŒ  Úname_cntÚ	var_namesÚarg_nameÚarg_sigÚvar_namerÞ  Úconst_tensorÚsymval_hintr‡  rŸ  Úcurrent_devicer¼   Ústream_names                       rz   Úcodegen_kernel_benchmarkz%TritonKernel.codegen_kernel_benchmark  s>  € ÜÓ!ˆØ,0¯I©I×,DÑ,DÓ,FÑ)ˆ�)˜Y¨à×ÑÒ5Ô6Ø�]‰]‹_ñ %	@Ü —‘Ó(ˆHØˆIÜ%(¨°IÓ%>ó  +Ñ!�˜'Ø!¤$ x£.Ð!1Ð2�Ü—g‘g×,Ñ,¨XÓ6�ÙØ×$Ñ$Ø#˜*Ð$4´Q·W±W×5EÑ5E×5PÑ5PÐQT×Q]ÑQ]ÓQ_Ó5`Ð4aÐacÔde×dkÑdk×dtÑdt×dÑdð  AD÷  AOñ  AOó  AQó  eRð  dSð  S]ð  ^a÷  ^lñ  ^ló  ^nð  ]oð  oxð  y|÷  yFñ  yFó  yHð  xIð  IJð  Köð ¤§¡×!2Ñ!2Ñ2ä#$§7¡7×#4Ñ#4°XÑ#>�LØ×$Ñ$Ø#˜*Ð$4´Q·W±W×5EÑ5E×5PÑ5PÐQ]×QbÑQbÓQdÓ5eÐ4fÐfhÔij×ipÑip×iyÑiy÷  jEñ  jEð  FR÷  FYñ  FYó  F[ó  j\ð  i]ð  ]gð  ht÷  h{ñ  h{ð  g|ð  |Eð  FR÷  FXñ  FXð  EYð  YZð  [öô   ¬Ô1Ü"#§'¡'×"2Ñ"2×"<Ñ"<¸W¿\¹\Ó"J�Kð
 %¨¯©Ñ4Ø&'˜Ø×$Ñ$¨ z°°[°MÐ%BÕCÜ ¬Ô6ÜŸW™W×@Ñ@ÓB�FÜŸG™G×,Ñ,×6Ñ6°w·}±}ÓE�EØ×$Ñ$Ø#˜* O°E°7¸*ÀVÀHÈIÐV]×VcÑVcÐUdÐdeÐfõô #ØDÀXÀJÐOóð ð × Ñ  Ö*ðA +ðB ×Ñ˜T×=Ñ=Ó?Ô@Ø×Ñ˜w t§y¡y°Ó';Ð&<¸AÐ>Ô?÷K%	@ðN 	×ÑÒ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Ø×ÑØLôð ×Ñ˜y¨¨Ð1Ô2Ø×ÑÐ=Ô>Ø×ÑØNô÷	ð  ˆ÷g%	@ñ %	@ú÷Xñ ú÷
	ñ 
	ú÷ ð ú÷	ð 	ú÷	ð  ˆúsf   ÁL2YÏAYÐ!BYÒ$YÓAY4Ô$A%Y(Ö	Y4Ö<A;Z ÙYÙY	ÙYÙY%Ù(Y1	Ù-Y4Ù4Y=Ú Z
c                ó˜   — t        j                  dj                  t        j                  j
                  j                  d«      «      «      S )Nzl
            from torch._dynamo.testing import rand_strided
            {}
            import torch
        Úget_raw_stream)ÚtextwrapÚdedentrD  rC   r  r‹  Úimport_get_raw_stream_asr¾   s    rz   Úimports_for_benchmark_kernelz)TritonKernel.imports_for_benchmark_kernel_  s:   € Ü�‰ð÷ ‰F”1—7‘7×%Ñ%×>Ñ>Ð?OÓPÓQó
ð 	
r|   c                óˆ   — | j                   ry| j                  ry| j                  r| j                  sJ ‚y| j                  ryy)Nr™  r©  rÏ  rë  Ú	pointwise)r™  r©  rÏ  r$  r¾   s    rz   Ú_get_heuristiczTritonKernel._get_heuristich  sD   € Ø×ÒØ!Ø×'Ò'Ø*Ø×&Ò&Ø×(Ò(Ð(Ð(Ø)Ø×"Ò"ØØr|   c                 óÔ  — t         j                  j                  j                  «       t        j                  «       t
        j                  t
        j                  t
        j                  j                  t
        j                  t
        j                  t
        j                  t
        j                  t
        j                  t
        j                  j                  t
        j                  j                   t
        j                  j"                  dœ} t         j$                  j&                  �d| d<   t        j(                  «       rd| d<   t
        j*                  rLt
        j*                  | d<   t
        j,                  | d<   t
        j.                  | d<   t
        j0                  | d<   t
        j2                  r9t
        j2                  | d	<   t
        j4                  | d
<   t
        j6                  | d<   | S )N)Úbackend_hashÚ$are_deterministic_algorithms_enabledÚ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_cubinTÚis_hiprw  Ú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)rs   ÚutilsÚ_tritonÚtriton_hash_with_backendr§  r   r¨  r©  rŒ   rª  r«  r¬  r­  r®  r¯  r°  r±  r²  rx  ry  rw  r´  rµ  r¶  r·  r¸  r¹  rº  )Úinductor_metas    rz   Úinductor_meta_commonz!TritonKernel.inductor_meta_commont  su  € ô "ŸK™K×/Ñ/×HÑHÓJÜ49×4^Ñ4^Ó4`Ü(.×(GÑ(GÜ$*×$?Ñ$?Ü"(§-¡-×"BÑ"BÜ%+×%AÑ%AÜ$*×$?Ñ$?Ü$*×$?Ñ$?Ü"×/Ñ/Ü&,×&CÑ&CÜ%+§]¡]×%HÑ%HÜ%Ÿ}™}×<Ñ<Ü!Ÿ=™=×4Ñ4ñ
ˆô �=‰=×ÑÐ(Ø&*ˆ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                óÖ  ‡ ‡!— t        «       }i }| j                  j                  «       D ]‚  \  }}t        |«      r| j                  sŒt
        j                  j                  j                  |«      }t        |t        t        j                  f«      sd}nt        t        |«      «      }|||<   Œ„ |€˜|j                  t        «       «       t
        j                  j!                  «       j"                  }|dk(  r|j                  d«       n|j                  d«       t$        j&                  r|j                  | j)                  «       «       | j*                  j-                  «       \  Š }	Š!}	t/        ‰!«      D ]Ÿ  \  }
}t        |t0        «      sŒt3        t        j4                  |j6                  «      }|t
        j                  j                  j8                  v sŒbt1        |j:                  t
        j                  j                  j8                  |   «      ‰!|
<   Œ¡ t=        t>           «       }| j@                  D �]  }|| j*                  jB                  v r(|jE                  | j*                  jB                  |   «       || j*                  jF                  v rj|t
        j                  jH                  vrN|| jH                  vr@|jE                  t3        tJ        | j*                  jF                  |   «      jL                  «       || j*                  jN                  v sŒß| j*                  jN                  |   }t        |tP        «      rJ ‚|jE                  |«       �Œ tS        ‰ ‰!«      D ]O  \  }}t        |tT        «      sŒ|jV                  tX        jZ                  k(  sŒ5|jE                  |j:                  «       ŒQ t]        |«      }| j_                  «       D ]Z  }t1        |j`                  › d�|jb                  «      }‰!je                  |«       ‰ je                  tg        |j:                  «      «       Œ\ ˆ ˆ!fd„}| jh                  D ]K  }|jj                  r| jl                  rŒ|jn                  €Œ) ||j`                  jq                  «       › d�«       ŒM | jr                  r |d«       tu        ‰!| jv                  ‰ ¬	«      }|ty        jz                  t
        j                  j!                  «       «      i d
œ}t
        j                  j|                  xs t
        j                  j~                  }| j�                  «       j‚                  t…        | j†                  «      t?        tˆ        jŠ                  «      ||| jŒ                  | jŽ                  | j�                  dœ| j“                  «       ¥}| jr                  r| jl                  |d<   d }t$        j&                  st$        j”                  r| j—                  «       dz  }||d<   t™        ‰!«      g|d<   t›        ‰!«      D ]  }d|d   ‰!|   j:                  <   Œ || _N        | jŸ                  «        | j                   D ]$  }|j£                  d«       |j                  |«       Œ& | j¤                  r2d| j§                  «       › d| j¤                  j$                  ›d|›d|›d�	}n | j                  r;| j¨                  j«                  «       }d| j§                  «       › d|›d|› d|›d|›d�}nYd}t­        |«      dk(  rt­        t¯        ‰!«      «      dk(  rd}nd}d| j§                  «       › d|›d|› d|›d|›d | j°                  › d�}|j                  |«       |j£                  d!|xs t?        tˆ        j²                  «      › d"djµ                  d#„ ‰ D «       «      › d$�«       |j·                  «       5  | j¹                  |«       | j*                  j»                  «       D ]  \  }}|j£                  |› d%|› �«       Œ |j                  | j¼                  «       d d d «       t$        j&                  r |j                  | j¿                  |«      «       |jÁ                  «       S # 1 sw Y   ŒIxY w)&Ni    Úcpuz"triton_helpers.set_driver_to_cpu()z"triton_helpers.set_driver_to_gpu()rÛ  c                ó†   •— t        «       r‰j                  t        | «      «       ‰j                  t        | d¬«      «       y )NT)Úis_constexpr)r>   rc  rI   rG   )r”  Úargdefsr%  s    €€rz   Úadd_constexpr_argz6TritonKernel.codegen_kernel.<locals>.add_constexpr_argý  s/   ø€ ä-Ô/Ø× Ñ ¤¨hÓ!7Ô8Ø�N‰Nœ7 8¸$Ô?Õ@r|   r›   r  )Ú
size_dtyperÄ  )r%  r‡  r†  )Ú	grid_typer¦  Úkernel_nameÚmutated_arg_namesÚoptimize_memr   Únum_loadÚnum_reductionrÏ  g    eÍÍAÚkernel_num_gbÚconfigsrE   r†  r‡   z$
                @triton_heuristics.z(
                    config=zI,
                    filename=__file__,
                    triton_meta=z$,
                    inductor_meta=z;
                )
                @triton.jit
            z!(
                    size_hints=z%,
                    reduction_hint=r   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  ó<   K  — | ]  }|j                  «       –— Œ y ­wrq   )Ú	full_namer¯  s     rz   rÍ   z.TritonKernel.codegen_kernel.<locals>.<genexpr>n  s   è ø€ ÒCcÐVWÀAÇKÁKÇMÑCcùrú  r  r"  )arM   r%  r  r9   r$  rC   r  rù   Úsymbolic_hintrb  r6  r®   r>  r-   r�   r”   rˆ  rµ  r   Úbenchmark_kernelr¡  rC  r  r\  rR   r	   r¯   r  Úinv_precomputed_replacementsrA  r   r´   Ú	mutationsÚinput_buffersra  rZ  Úremoved_buffersrN   Ú
inner_nameÚoutput_buffersrQ   rý   rT   Ú	zero_moderU   ÚZERO_ON_CALLr»  rû  r!  rÛ  rc  rG   r,  r'  rÏ  Ú
tensor_dimr°   r©  r`   rr  r)   r6  Úis_inferenceÚis_backwardÚ_get_grid_typerr   Úsetr¦  r8   ÚDESCRIPTIVE_NAMEr   rË  rÌ  r¿  r´  Úestimate_kernel_num_bytesr\   r]   r§  rt  r•  r�   r™  r¤  r&  Úget_reduction_hintr  r^   r8  rŽ  rØ   rf  Úcodegen_static_numelsÚaliasesr›  r›  r‘   )"r¿   rA  Úcoder‚  r!  rÛ  Ú
numel_hintrw  Údevice_typerŒ  rŠ  r	  rÚ  Úmutated_argsÚmutationÚmutation_argÚargnamer¢   ÚsizeargrÅ  Útriton_meta_signaturer§  rÊ  r¾  r�  Úarg_numr  Úheuristics_lineÚreduction_hintÚ	tile_hintÚoldÚnewrÄ  r%  s"                                   @@rz   Úcodegen_kernelzTritonKernel.codegen_kernelœ  s|  ù€ ÜÓˆàˆ
Ø!Ÿ[™[×.Ñ.Ó0ò 	+‰MˆF�EÜ" 6Ô*°4×3HÒ3HØäŸ™×)Ñ)×7Ñ7¸Ó>ˆJÜ˜j¬3´·±Ð*>Ô?ð !‘	ä+¬C°
«OÓ<�	Ø!*ˆJ�vÒð)	+ð, ˆ<Ø�K‰KÔ1Ó3Ô4ÜŸ'™'×=Ñ=Ó?×DÑDˆKØ˜eÒ#Ø—‘Ð@ÕAà—‘Ð@ÔAä×&Ò&Ø—‘˜D×=Ñ=Ó?Ô@à#'§9¡9×#;Ñ#;Ó#=Ñ ˆ��I˜qä 	Ó*ò 	‰FˆAˆsÜ˜#œwÕ'ô œeŸl™l¨C¯H©HÓ5�ØœQŸW™W×-Ñ-×JÑJÒJÜ#*ØŸ™¤!§'¡'×"2Ñ"2×"OÑ"OÐPVÑ"Wó$�I˜a’Lð	ô "¤#‘Ó(ˆØŸ™ó 	/ˆHØ˜4Ÿ9™9×2Ñ2Ñ2Ø× Ñ  §¡×!8Ñ!8¸Ñ!BÔCà˜DŸI™I×5Ñ5Ñ5Ø¤A§G¡G×$;Ñ$;Ñ;Ø D×$8Ñ$8Ñ8à× Ñ Üœ¨¯©×)BÑ)BÀ8Ñ)LÓM×XÑXôð ˜4Ÿ9™9×3Ñ3Ò3Ø#Ÿy™y×7Ñ7¸ÑA�Ü% l´JÔ?Ð?Ð?Ø× Ñ  Ö.ð	/ô6   ¨Ó3ò 	/‰LˆG�Sä˜3¤Õ-Ø—M‘MÔ%6×%CÑ%CÓCà× Ñ  §¡Õ.ð	/ô ˜lÓ+ˆà×+Ñ+Ó-ò 	2ˆDÜ §¡ ¨UÐ3°T·Z±ZÓ@ˆGØ×Ñ˜WÔ%Ø�N‰Nœ7 7§<¡<Ó0Õ1ð	2õ	Að ×$Ñ$ò 	=ˆDØ× Ò  T×%>Ò%>àØ�‰Ð&Øá §¡×!2Ñ!2Ó!4Ð 5°UÐ;Õ<ð	=ð ×%Ò%Ù˜hÔ'ä 1Ø $×"2Ñ"2¸Gô!
Ðð /Ü&×-Ñ-¬a¯g©g×.QÑ.QÓ.SÓTØñ'
ˆô —w‘w×+Ñ+ÒB¬q¯w©w×/BÑ/Bˆð ×,Ñ,Ó.×7Ñ7Ü! $×"5Ñ"5Ó6Üœ{×;Ñ;Ó<Ø!-Ø(ØŸ™ØŸ™Ø!×/Ñ/ñ
ð ×'Ñ'Ó)ð
ˆð ×%Ò%Ø48×4MÑ4MˆMÐ0Ñ1àˆÜ×"Ò"¤f×&>Ò&>Ø×3Ñ3Ó5¸Ñ;ˆFØ-3ˆM˜/Ñ*ä"+¨IÓ"6Ð!7ˆ�IÑô +¨9Ó5ò 	BˆGØ@AˆK˜Ñ$ Y¨wÑ%7×%<Ñ%<Ò=ð	Bð 'ˆÔà×ÑÔà×+Ñ+ò 	 ˆFØ�N‰N˜2ÔØ�K‰K˜Õð	 ð ×Òð#$Ø$(×$7Ñ$7Ó$9Ð#:ð ;Ø ×-Ñ-×4Ñ4Ð7ð 8!à!, ð 0#Ø#0Ð"3ð 4ð‰Oð ×"Ò"Ø!Ÿ]™]×=Ñ=Ó?ˆ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Ô$Ø�‰Ø�4Ò7œ3œ{×6Ñ6Ó7Ð8¸¸$¿)¹)ÑCcÐ[bÔCcÓ:cÐ9dÐdfÐgô	
ð �[‰[‹]ñ 	#Ø×&Ñ& tÔ,Ø ŸI™I×-Ñ-Ó/ò 1‘��SØ—‘ #  c¨#¨Ð/Õ0ð1à�K‰K˜Ÿ	™	Ô"÷		#ô ×"Ò"Ø�K‰K˜×5Ñ5°fÓ=Ô>à�}‰}‹Ð÷	#ð 	#ús   Þ2A%aáa(c                ó¦  — t         j                  j                  j                  | «      } t	        | t
        j                  t        f«      rt        | «      }t        |«      }|S d}t         j                  j                  j                  | |«      sC|dkD  rt        d| › �«      ‚|dz  }t         j                  j                  j                  | |«      sŒC|S )Né€   i @  z!Failed to find static RBLOCK for r   )rC   r  rù   Úsimplifyrb  r®   r>  r6  r-   Ústatically_known_leqrà  )r[  r3  s     rz   Ú_get_persistent_RBLOCKz#TritonKernel._get_persistent_RBLOCK{  s°   € ä—‘×!Ñ!×*Ñ*¨6Ó2ˆÜ�fœuŸ}™}¬cÐ2Ô3Ü�f“+ˆCÜ! #Ó&ˆCð ˆ
ð ˆCÜ—g‘g×&Ñ&×;Ñ;¸FÀCÔHØ˜’?Ü$Ð'HÈÈÐ%QÓRÐRØ�q‘�ô —g‘g×&Ñ&×;Ñ;¸FÀCÕHð ˆ
r|   c                óN   — 	 t         j                  | «       y# t        $ r Y yw xY w)NTF)r”  rù  rà  )r[  s    rz   Úhas_persistent_RBLOCKz"TritonKernel.has_persistent_RBLOCK‰  s*   € ð	Ü×/Ñ/°Ô7ØøÜò 	Ùð	ús   ‚ ˜	$£$c                óÆ  — d	d„}| j                   D �]M  }|j                  r| j                  rdt        j                  j
                  j                  |j                  «      } ||«      r)|j                  |j                  › dt        |«      › �«       |j                  r’| j                  r†| j                  r1| j                  | j                  |j                  «      «      }d|› d�}n| j                  |j                  «      }|j                  |j                  j!                  «       › d|› �«       |j                  dk(  s�Œ/| j"                  s�Œ=|j                  d«       �ŒP y)
aÑ  
        We get a small speedup from hard coding numels if they are static.

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

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

        We would add
        xnumel = 4096
        r0_numel = 768

        After the signature, before the kernel code, if we decided to make these static. As its hardcoded, it becomes
        a better signal to triton on how to unroll and do some static indexing. So, it's not so much that downstream
        knows that its a static numel, as that you just plop a constant into the kernel.
        c                óB   — t        | t        j                  t        f«      S rq   )rb  r®   r>  r6  )r  s    rz   Úis_static_integerz=TritonKernel.codegen_static_numels.<locals>.is_static_integer£  s   € Ü˜d¤U§]¡]´CÐ$8Ó9Ð9r|   znumel = z*triton_helpers.constexpr_next_power_of_2((z + RSPLIT - 1) // RSPLIT)zBLOCK: tl.constexpr = r  zXBLOCK: tl.constexpr = 1N)r  r»   r~   r}   )r,  r'  r$  rC   r  rù   r÷  rÛ  r�   r!  r6  rÏ  r©  r–  r3  rù  r°   r   )r¿   rå  rþ  r¢   Úsimplified_tree_numelrÛ  r3  s          rz   rã  z"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Ø—‘ $§+¡+×"3Ñ"3Ó"5Ð!6Ð6LÈSÈEÐRÔSà�{‰{˜cÔ! d§m¤mØ—‘Ð9Ö:ñ	;r|   c                óÌ  — t        | j                  D �cg c]  }t        |j                   «      ‘Œ c}«      }| j                  r|dk(  sJ ‚t
        j                  S |dk(  rt
        j                  S |dk(  rIt        t        | j                  | j                  «      «      rt
        j                  S t
        j                  S |dk(  rt
        j                  S t        d|› �«      ‚c c}w )NrE   r   r   z"Unsupported number of dimensions: )rþ  r,  r6  r'  r©  r&   ÚCooperativeReductionGridÚGrid1DrÐ   rÙ   Úneeds_yz_grid_overflowÚGrid2DWithYZOverflowÚGrid2DÚGrid3Drà  )r¿   r¢   r  s      rz   rÞ  zTritonKernel._get_grid_type·  sÄ   € Ü¸×8HÑ8HÖI°”˜×*Ñ*Ð*Õ+ÒIÓJˆØ×%Ò%Ø˜’6ˆM�6Ü$×=Ñ=Ð=Ø�!ŠVÜ$×+Ñ+Ð+Ø�!ŠVÜ”3�t×2Ñ2°D×4DÑ4DÓEÔFÜ(×=Ñ=Ð=Ü$×+Ñ+Ð+Ø�!ŠVÜ$×+Ñ+Ð+ÜÐ=¸a¸SÐAÓBÐBùò Js   ”C!c                ó„  — | j                   D ]±  }t        |j                  t        j                  t        j
                  f«      r|j                  }n*t        j                  j                  j                  ||«      }|j                  r| j                  sŒ‡|j                  |«       |j                  t        |«      «       Œ³ y rq   )r,  rb  rÛ  r®   r>  r¯   rC   r  Úwrapper_codeÚgenerate_numel_exprr'  r$  rc  rµ  )r¿   rA  r‘  Ú	arg_typesr¢   r  s         rz   rv  z#TritonKernel.add_numel_to_call_argsÆ  sˆ   € à×$Ñ$ò 	-ˆDÜ˜$Ÿ*™*¤u§}¡}´e·l±lÐ&CÔDØ—z‘z‘ä—w‘w×+Ñ+×?Ñ?ÀÀdÓK�à×$Ò$¨×(=Ó(=Ø× Ñ  Ô&Ø× Ñ ¤ d£Õ,ñ	-r|   c                ó¾  — t         j                  j                  }|j                  «        | j                  j                  «       \  }}}}| j                  |||«       | j                  j                  D ]  }|j                  |«       Œ |j                  ||d|| j                  ¬«       t        | j                  j                  «      D ]  }|j                  |«       Œ y )NT)rŒ   r
  r§  )rC   r  r  Úwrite_triton_header_oncerC  r  rv  Úworkspace_argsÚgenerate_workspace_allocationÚgenerate_kernel_callr§  r)  Úgenerate_workspace_deallocation)r¿   rA  rC  ÚwrapperrŒ  r‘  r
  Úwss           rz   Úcall_kernelzTritonKernel.call_kernelÒ  sÉ   € Ü—'‘'×&Ñ&ˆØ×(Ñ(Ô*Ø%)§Y¡Y×%=Ñ%=Ó%?Ñ"ˆˆ9�a˜Ø×#Ñ# D¨)°YÔ?à—)‘)×*Ñ*ò 	6ˆBØ×1Ñ1°"Õ5ð	6ð 	×$Ñ$ØØØØØ×(Ñ(ð 	%ô 	
ô ˜4Ÿ9™9×3Ñ3Ó4ò 	8ˆBØ×3Ñ3°BÕ7ñ	8r|   c                ó‚  — t         j                  j                  }| j                  j	                  «       \  }}}}t        ||«      D ]w  \  }}t        |t        «      sŒt         j                  j                  r|j                  d|› d|› d�«       ŒJd|› d�}|j                  |«       d|› d�}|j                  |«       Œy y )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())
rC   r  r  rC  r  rý   rb  rS   Úcpp_wrapperr�   )r¿   r  rŒ  r‘  Úarg_signaturesr	  Úarg_signaturer8  s           rz   Úcodegen_nan_checkzTritonKernel.codegen_nan_checkæ  sÀ   € Ü—'‘'×&Ñ&ˆØ*.¯)©)×*BÑ*BÓ*DÑ'ˆˆ9�n aÜ"% i°Ó"@ò 
	,ÑˆC�Ü˜-¬Õ3Ü—7‘7×&Ò&Ø×%Ñ%ØTÐUXÐTYÐY\Ð]`Ð\aÐadÐeõð )¨¨Ð-BÐC�DØ×%Ñ% dÔ+Ø(¨¨Ð-BÐC�DØ×%Ñ% dÕ+ñ
	,r|   c                ó   — t        |i |¤ŽS rq   )r   )r¿   rC  r  s      rz   Úcreate_cse_varzTritonKernel.create_cse_varõ  s   € Ü  $Ð1¨&Ñ1Ð1r|   c                ó  — |j                   › d| j                  | j                  |j                  «      «      › �}|j                  j
                  r| j                  j                  |«       y | j                  j                  |«       y r   )	rA  r–  r3  r  ÚrootrÀ  rd  r�   r›  )r¿   Úentryr8  s      rz   Úcodegen_iteration_ranges_entryz+TritonKernel.codegen_iteration_ranges_entryø  sd   € Ø—*‘*�˜S §¡¨D×,@Ñ,@ÀÇÁÓ,LÓ!MÐ NÐOˆØ�:‰:×ÒØ×Ñ×(Ñ(¨Õ.ð �I‰I×Ñ Õ%r|   c                ó  — |j                   €J ‚| j                  |j                   «      }| j                  }|dk7  rd|› d�nd}| j                  r| j                  r|j
                  r|› d�}d|j                  j                  «       › d|› |› �S )Nr'  r  rü   r‡   z + rsplit_startztl.arange(0, zBLOCK))rÛ  Úindexing_size_strrr  r©  rÏ  r'  r!  r°   )r¿   r  rj  rr  r›  s        rz   rÄ  z)TritonKernel.iteration_ranges_ranges_code   s˜   € Ø×ÑÐ+Ð+Ð+Ø×%Ñ% e×&6Ñ&6Ó7ˆØ×&Ñ&ˆØ*5¸Ò*C�4˜�} AÑ&Èˆà×&Ò&Ø×)Ò)Ø×"Ò"à�x˜Ð/ˆFØ˜uŸ|™|×1Ñ1Ó3Ð4°F¸4¸&ÀÀÐIÐIr|   c                ó^   — | j                   }| j                  «       }dg|z  }d|› d|› d|› d�S )NrE   rJ  rû   rü   )rr  r"  )r¿   r  r  rr  r#  rj  s         rz   Úiteration_ranges_scalar_codez)TritonKernel.iteration_ranges_scalar_code  sC   € ð ×&Ñ&ˆØ×&Ñ&Ó(ˆØˆs�T‰zˆØ˜$˜˜r % ¨¨;¨-°qÐ9Ð9r|   c                ó$  — |j                   €J ‚d|j                   › d�}| j                  |«      r#d|› d|j                   dz   › d|j                   › d�}|j                  j                  ||«      }| j                  dk7  r|› d	| j                  › d�S |S )
Nztl.program_id(rü   r®  z + tl.program_id(rE   z) * tl.num_programs(r´  r'  r  )r¶  r  Ú	pid_cacher^  rr  )r¿   r  r2  Úpids       rz   Úiteration_ranges_get_pidz%TritonKernel.iteration_ranges_get_pid  s©   € Ø�~‰~Ð)Ð)Ð)Ø˜uŸ~™~Ð.¨aÐ0ˆð ×&Ñ& uÔ-ð �c�UÐ+¨E¯N©N¸QÑ,>Ð+?Ð?SÐTY×TbÑTbÐScÐceÐfˆCØ�o‰o×!Ñ! # sÓ+ˆØ×Ñ˜zÒ)Ø�U˜$˜t×/Ñ/Ð0°Ð2Ð2Øˆ
r|   c                óÚ   — |j                   dk(  xr[ |j                   xrL | j                   xr= t        j                  j
                  j                  |j                  t        «       «       S r‚  )	r¶  Úhas_zdimr©  rC   r  rù   rø  rÛ  r,   )r¿   r  s     rz   r  z#TritonKernel.needs_yz_grid_overflow$  sa   € à�N‰N˜aÑò YØ—N‘NÐ"òYà×.Ñ.Ð.òYô —G‘G×$Ñ$×9Ñ9¸%¿+¹+Ä~ÓGWÓXÐXð		
r|   c                óˆ   — | j                   r | j                   |j                  «       › d�   S t        |j                  «          S )Nr›   )r™  r°   r*   )r¿   r!  s     rz   râ  zTritonKernel.max_block,  s;   € Ø×ÒØ×$Ñ$¨¯©«Ð'7°uÐ%=Ñ>Ð>Ü §¡£Ñ/Ð/r|   c                óœ  — | j                   sy| j                  rW|j                  j                  «       › d�| j                  v r.| j                  |j                  j                  «       › d�   dk(  r6yt        j
                  j                  j                  |j                  d«      ry|j                  r(| j                  r| j                  |j                  «      }n9|j                  dk(  r| j                  rd}n| j                  |j                  «      }|j                  r| j                  r|| j                  «       z  }t        j
                  j                  j!                  |j                  |«      r[|j"                  dk7  xsJ |j$                  xs< t        j
                  j                  j'                  |j                  t)        «       «      S y)NFr›   rE   Tr  )r˜  r™  r!  r°   rC   r  rù   rö   rÛ  r'  rÏ  rù  r   râ  r©  r  rG  r¶  r(  rø  r,   )r¿   r¢   râ  s      rz   rµ  zTritonKernel._has_constant_mask1  sa  € Ø×!Ò!Øà×Ò D§K¡K×$5Ñ$5Ó$7Ð#8¸Ð!>À$×BSÑBSÑ!SØ× Ñ  D§K¡K×$5Ñ$5Ó$7Ð#8¸Ð!>Ñ?À1ÒDØä�w‰w×Ñ×7Ñ7¸¿
¹
ÀAÔFØð ×Ò ×!:Ò!:Ø×3Ñ3°D·J±JÓ?‰IØ�[‰[˜CÒ D§M¢MØ‰IàŸ™ t§{¡{Ó3ˆIà×Ò ×!;Ò!;Ø! D§O¡OÓ$5Ñ5ˆIô �7‰7×Ñ×8Ñ8¸¿¹ÀYÔOà—‘ Ñ"ò WØ—=‘=òWä—7‘7×#Ñ#×8Ñ8¸¿¹Ä^ÓEUÓVðð r|   c                ód   — | j                   d   }|j                  dk(  sJ ‚| j                  |«      S )Nr   r  )r,  r!  rµ  )r¿   Úxtrees     rz   r¼  z TritonKernel._has_constant_xmask[  s5   € Ø× Ñ  Ñ#ˆØ�|‰|˜sÒ"Ð"Ð"Ø×&Ñ& uÓ-Ð-r|   c                ó¨   — | j                   D ]2  }| j                  |«      sŒ|j                  |j                  › d�«       Œ4 |j                  d«       y )NrÌ   r   )r,  rµ  r@  r!  )r¿   r·   r¢   s      rz   rÿ  zTritonKernel.filter_masks`  sN   € Ø×$Ñ$ò 	8ˆDØ×&Ñ& tÕ,Ø×!Ñ! T§[¡[ M°Ð"6Õ7ð	8ð
 	×Ñ˜&Õ!r|   c                ó‚   — t        t        j                  «      d | j                   D �cg c]  }t        |   ‘Œ c}S c c}w rq   )r(  r–   r©   r#  r   )r¿   r    s     rz   Úget_reduction_prefixesz#TritonKernel.get_reduction_prefixesh  sB   € ô œ]×:Ñ:Ó;Ð<U¸d×>UÑ>UÐVö
àô �tÓò
ð 	
ùò 
s   ©<c                óÈ  — | j                   D �cg c]  }|j                  sŒ|‘Œ }}dj                  t        d„ |D «       «      «      }|j	                  d| j                  |«      › �«       | j                   D �cg c]+  }|j                  rt        j                  |j                     ‘Œ- }}t        |«      }|j	                  d| j                  |«      › �«       yc c}w c c}w )z^
        Generates code that flattens ND reduction numels, block sizes, etc. into 1D.
        rß  c              3  ó:   K  — | ]  }|j                   › d �–— Œ y­w)rÛ  Nr†  rÁ  s     rz   rÍ   z8TritonKernel.codegen_reduction_numels.<locals>.<genexpr>u  s   è ø€ Ò"U¸T d§k¡k ]°%Ô#8Ñ"Uùrˆ  z	rnumel = zRBLOCK: tl.constexpr = N)
r,  r'  rØ   r»  r�   r–  r–   rŸ   r    r;   )r¿   r9  r¢   Úreduction_treesr[  Ú	rn_blocksrN  s          rz   r¨  z%TritonKernel.codegen_reduction_numelso  sÍ   € ð
 -1×,<Ñ,<ÖR DÀ×@QÓ@Qš4ÐRˆÐRØ—‘œFÑ"UÀ_Ô"UÓUÓVˆØ�‰˜	 $§*¡*¨VÓ"4Ð!5Ð6Ô7ð
 ×(Ñ(ö
àØ× Ò ô ×%Ñ% d§i¡iÓ0ð
ˆ	ð 
ô
 ˜yÓ)ˆØ�‰Ð/°·
±
¸6Ó0BÐ/CÐDÕEùò Sùò

s   �C¡CÁ90Cc                ó|   — | j                  «       }|D �cg c]  }t        j                  |› |› �fi |¤Ž‘Œ c}S c c}w )zK
        Helper to initialize symbols like rn_numel, rn_base, etc.
        )r/  r®   r¯   )r¿   r›  r  Úrn_prefixesr!  s        rz   rÅ  z#TritonKernel._get_reduction_symbols�  s=   € ð ×1Ñ1Ó3ˆØJUÖVÀ”—‘ ˜x¨ xÐ0Ñ;°FÓ;ÒVÐVùÒVs   •!9c                óê   — | j                  «       }| j                  ddd¬«      }t        t        |«      dz
  «      D �cg c]  }t	        ||dz   d «      ‘Œ c}t        j                  d«      gz   S c c}w )z¸
        Compute coefficients to convert ND reduction indices to linear indices.
        For example:
          rindex = r0_index * r1_numel * ... * rn_numel + ... + rn_index.
        rÛ  Trœ   rE   N)r/  rÅ  r*  r  r;   r®   r>  )r¿   r5  Ú	rn_numelsrJ  s       rz   Ú_get_reduction_index_coeffsz(TritonKernel._get_reduction_index_coeffsˆ  s{   € ð ×1Ñ1Ó3ˆØ×/Ñ/°ÀÐPTÐ/ÓUˆ	ä;@ÄÀ[ÓAQÐTUÑAUÓ;Vö
Ø47ŒM˜) C¨!¡G IÐ.Õ/ò
ä�]‰]˜1ÓÐñð 	ùò 
s   ¾A0c                ó:   — | j                  «       }t        ||«      S )zK
        Compute linear reduction indices from N dimensional ones.
        )r8  r:   )r¿   Ú
multi_indsÚcoeffss      rz   rÆ  z'TritonKernel._flatten_reduction_indices•  s   € ð ×1Ñ1Ó3ˆÜ˜ Ó,Ð,r|   c                ó$  — | j                  ddd¬«      }| j                  ddd¬«      }| j                  |«      }|j                  d| j                  |«      › �«       | j                  |«      }|j                  d| j                  |«      › �«       y)zX
        Generates code that converts ND reduction indices into linear indices.
        r—   Tr˜   r¼   z
roffset = z	rindex = N)rÅ  rÆ  r�   r  )r¿   r9  Ú
rn_offsetsÚrn_indsr9  Úrindexs         rz   rÇ  z&TritonKernel.codegen_reduction_indicesœ  sš   € ð
 ×0Ñ0Ø˜d°ð 1ó 
ˆ
ð ×-Ñ-¨g¸tÐQUÐ-ÓVˆð ×1Ñ1°*Ó=ˆØ�‰˜
 4×#4Ñ#4°WÓ#=Ð">Ð?Ô@Ø×0Ñ0°Ó9ˆØ�‰˜	 $×"3Ñ"3°FÓ";Ð!<Ð=Õ>r|   c                óØ  — |j                   }|j                  r%|j                  |j                  › d|› d|› d�«       nÐ|j                  €D|j                  |j                  › d| j                  |«      › �«       |j                  |› d�«       n€|j                  �|› d| j                  |«      › �}n| j                  ||› d�«      }|j                  |› d| j                  |«      › d|j                  «       › d�|j                  › d|› �g«       | j                  |«      r(| j                  «       }|j                  |› d	|› d
�«       y |j                  |› d|j                  › d|› d�«       y )Nr"  z	offset + rÂ  z
offset = 0r—   z	offset = rß  r›   zmask = tl.full(z, True, tl.int1)zmask = z < rÛ  )r!  rÀ  r�   rA  r¶  rÄ  rÛ  r"  r€  r&  r°   rµ  r  )r¿   r  rå  r  r8  rx  s         rz   rÃ  z,TritonKernel.iteration_ranges_codegen_header¬  sb  € ð �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¸À1À#ÀVÀÓM�Ø�O‰Oà�c˜ 4×#@Ñ#@ÀÓ#GÐ"HÈÈAÏGÉGËIÈ;ÐV[Ð\Ø—z‘z�l # d VÐ,ðôð ×"Ñ" 5Ô)Ø×'Ñ'Ó)ˆEØ�N‰N˜a˜S °¨wÐ6FÐGÕHà�N‰N˜a˜S ¨¯
©
 |°3°q°c¸Ð?Õ@r|   )r   TN)r­  zdict[str, sympy.Expr]r™  zOptional[FixedTritonConfig]r~   r   ©rè  r  r~   r´   rÛ   rÜ   )r¼   r»   )r‡   )rA  r´   rý  r´   r)  rß   r~   ztuple[str, str])r  r»   rj  r»   r6  r}   r°   r}   )rA  r´   r¼   r»   rq   )
rA  r´   r¼   r»   r  rK   ra  rB   r~   r   ©NN)r!  rK   ri  z.tuple[str, sympy.Expr, sympy.Expr, sympy.Expr]rj  rK   rk  r  r  r}   rl  z Optional[tuple[str, sympy.Expr]]rm  zOptional[CSEVariable]r~   rK   )r  r´   rè  r  r~   r´   )
rè  r  r2  r  r”  rA   r  ú+Union[CSEVariable, tuple[CSEVariable, ...]]r~   rC  )rè  r  )rA  r´   r¼   r»   r  rC  )r  útuple[torch.dtype, ...]r~   r´   )r  rD  rÛ  zUCallable[[tuple[CSEVariable, ...], tuple[CSEVariable, ...]], tuple[CSEVariable, ...]]r!  útuple[CSEVariable, ...]r~   rE  )
r  rD  r!  rE  r]  r}   r^  r}   r~   rE  )r~   rW  )r~   z type[triton_heuristics.GridExpr])rA  r´   rC  zOptional[IRNode]rg  )r~   r   )r  rX   )r  rY   r~   r´   )r  rY   r  r   r~   r´   )r  rY   r~   r}   )r!  r´   r~   r6  )r¢   rY   r~   r}   )r·   r¶   r~   r   )r9  rM   r~   r   )r›  r´   r~   zlist[sympy.Symbol]rY  )r:  r  r~   r»   )r  rY   rå  rM   r~   r   )Mrr   r€   r�   r   r  rƒ   r2  r–  r  r  r¯  r²  rª  r¬  r«  rË  rÍ  rÐ  rÝ   rÒ  r)  r)  r-  r:  r5  r9  rd  r_  rt  ry  r  rë  rï  rÁ  r«  rÊ  rÇ  r  rÉ  r  r'  rS  r`  rt  rz  r›  r¡  r¤  r\  r¿  rô  rù  rû  rã  rÞ  rv  r  r  r  r  rÄ  r"  r&  r  râ  rµ  r¼  rÿ  r2   r/  r¨  rÅ  r8  rÆ  rÇ  rÃ  r  r  s   @rz   r”  r”  1  sÖ  ø… Ø%€IØ%Ó%Ø).€EÐ&Ó.Ø€Oð
 ØØ48ð%3à%ð%3ð
 2ð%3ð 
õ%3óN"ó
ò
#òJò*:ò0ó
ò
ð ò"ó ð"ð ØØØñfTàófTðR	 EGð- Øð- Ø!ð- Ø-<ð- à	ó- ó^8ðMàðMð ðMð ð	Mð
 óMò4óEðP SWð(Øð(Ø *ð(Ø3>ð(ØFOð(à	ó(òTð  48Ø04ð1àð1ð Cð1ð &ð	1ð
 $ð1ð ð1ð 1ð1ð .ð1ð 
ó1óf.ó
ð"DàðDð ðDð &ð	Dð
 ;ðDð 
5óDðL&ØDOó&ó"ò9
òv5ò(&ò!ò
%ð6(àð(ð ð(ð ;ó	(óT1Sðfm"à'ðm"ð
ðm"ð (ðm"ð 
!óm"ð^>"à'ð>"ð (ð>"ð ð	>"ð
 ð>"ð 
!ó>"ò@b%óHò Xòt
ò
ð ñ%ó ð%óN]ð~ ñó ðð ñó ðò$;óLCò
-ô8ó(,ó2ó&óJð:Ø(ð:Ø14ð:à	ó:óó
ó0ó
(óT.ó
"ð ò
ó ð
óFó$Wð ò
ó ð
ó-ó?ð AØ(ðAØ0>ðAà	÷Ar|   r”  c            
      ó`  ‡ — e Zd ZU eZded<    eej                  ej                  ej                  ej                  ej                  ej                  ej                  ej                  g«      Zdˆ fd„Zedd„«       Zd„ Zd„ Zddd„Z	 d	 	 	 dd„Z	 	 	 	 	 	 	 	 dd	„Z	 	 	 	 	 	 	 	 dd
„Zd„ Zˆ xZS )ÚTritonSchedulingz	type[Any]Úkernel_typec                ó¨   •— t         ‰| �  |«       |�t        |d«      sy |j                  D ]$  }t	        |t
        t        f«      sŒt        |_        Œ& y )NrA  )	r  r  rŠ   rA  rb  r1   r/   Údebug_triton_codeÚdebug_device_str)r¿   Ú	schedulerrC  r  s      €rz   r  zTritonScheduling.__init__Ø  sM   ø€ Ü‰Ñ˜Ô#ØÐ¤G¨I°wÔ$?ØØ—O‘Oò 	:ˆDÜ˜$¤Ô0BÐ CÕDÜ(9�Õ%ñ	:r|   c                óÐ   — t         j                  j                  st         j                  j                  r't	        g | j
                  ¢t        j                  ‘«      S | j
                  S rq   )r   rŒ   Úcooperative_reductionsÚforce_cooperative_reductionsr   Úbackend_featuresrH   ÚREDUCE_TO_SINGLE_ELEMENT)rv   r‡  s     rz   Úget_backend_featuresz%TritonScheduling.get_backend_featuresà  sR   € ô �M‰M×0Ò0Ü�}‰}×9Ò9äØP�#×&Ñ&ÐP¬×(OÑ(OÐPóð ð ×#Ñ#Ð#r|   c                ó”  ‡— t         j                  j                  }t        ||«      \  }}|r|j	                  |«       t
        j                  rvddlm}m	Š t        ˆfd„|D «       «      sY|D �cg c]  }t        ||«      r|j                  «       ‘Œ  }}|j	                  |j                  › ddj                  |«      › �«       y y y c c}w )Nr   )r.   ÚForeachKernelSchedulerNodec              3  ó6   •K  — | ]  }t        |‰«      –— Œ y ­wrq   )rb  )r±   r  rT  s     €rz   rÍ   z3TritonScheduling.codegen_comment.<locals>.<genexpr>÷  s   øè ø€ ò Ø>?”
˜1Ð8×9ñùs   ƒz Fused node name list: rû   )rC   r  r  r6   r�   r   Údebug_fusionÚtorch._inductor.schedulerr.   rT  rÐ   rb  Úget_nameÚcommentrØ   )	r¿   Únode_scheduler  ÚoriginsÚ_detailed_originsr.   r  Ú
node_namesrT  s	           @rz   Úcodegen_commentz TritonScheduling.codegen_commentë  sÇ   ø€ Ü—'‘'×&Ñ&ˆÜ%8¸ÈÓ%PÑ"ˆÐ"ÙØ×Ñ˜gÔ&ä×Ò÷ô
 ó ØCPôô ð +öàÜ! !Ð%6Ô7ð —J‘J•Lð�
ð ð
 ×!Ñ!Ø—‘Ð'Ð'>¸t¿y¹yÈÓ?TÐ>UÐVõðð ùòs   Á.#Cc                óÚ  — t         j                  j                  }||j                  v r|j                  |   }|S t        j
                  j                  r$t        |t        j
                  j                  «      nd}t        |«      d d }dj                  d|||j                  «       g«      }||j                  |<   t        j
                  j                  r|nd}|j                  t        t        j                  «      |«      }|j                  t        t        j                   «      |«      }|j                  dd«      }t#        t%        |j'                  «       «      d«      \  }	}
}t)        «       }t*        j-                  «       rt*        j                  ||«       |j/                  d	|›d
�«       |j1                  |d¬«       t         j                  j3                  «       }|j/                  d|j4                  › d�«       d|› �}t7        ||«      \  }}|d|z   dz   |z   z  }|j9                  ||j;                  «       |«       t=        j>                  d«      rt=        j@                  |||«       |S )Nr‡   r   rŒ  rŒ   Útriton_z#pragma CMTú#Úpyzasync_compile.triton(z, '''Tr  z''', device_str='z')z# kernel path: r}  Úkernel_metadata)!rC   r  r  Úsrc_to_kernelr   rŒ   Údescriptive_namesr5   rD   rØ   Únext_kernel_suffixÚunique_kernel_namesÚreplacer´   r8   rà  rŽ  r#   r"   r  rM   Úasync_compileÚuse_process_poolr�   r�   rˆ  rµ  r6   Údefine_kernelr‘   r    Úis_metric_table_enabledÚlog_kernel_metadata)r¿   Úsrc_coderZ  r  r  rÈ  Ú
fused_nameÚkernel_categoryÚ	subs_nameÚ	_basenamerŒ  Úkernel_pathÚcompile_wrapperr™  Úmetadata_commentr[  Údetailed_originss                    rz   rk  zTritonScheduling.define_kernel  s4  € Ü—'‘'×&Ñ&ˆØ�w×,Ñ,Ñ,Ø!×/Ñ/°Ñ9ˆKðj Ðôc —=‘=×2Ò2ô & m´V·]±]×5TÑ5TÔUàð ô
 AÀÓJÈ2ÈAÐNˆOØŸ(™(Ø˜?¨J¸×8RÑ8RÓ8TÐUóˆKð /:ˆG×!Ñ! (Ñ+Ü'-§}¡}×'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�q˜+Ü,Ó.ˆOä×-Ñ-Ô/ô ×$Ñ$ Y°Ô9à×%Ñ%Ð(=¸i¸]È%Ð&PÔQà×"Ñ" 8°4Ð"Ô8ÜŸW™W×@Ñ@ÓBˆNØ×%Ñ%Ð(9¸.×:MÑ:MÐ9NÈbÐ&QÔRà!0°°Ð>ÐÜ(;¸MÈ7Ó(SÑ%ˆGÐ%Ø  w¡°Ñ 5Ð8HÑ HÑHÐØ×!Ñ!Ø˜_×5Ñ5Ó7Ð9Iôô ×.Ñ.Ð/@ÔAÜ×+Ñ+¨K¸ÀhÔOàÐr|   c                óš   — | j                  |d¬«      }t        j                  |«      }| j                  ||t	        d„ |D «       «      ¬«      S )zƒ
        Benchmark fused list of nodes and return the execution time
        in milliseconds on randomly generated inputs.
        T)rÒ  c              3  ó<   K  — | ]  }|j                  «       –— Œ y ­wrq   ©rX  ©r±   r  s     rz   rÍ   z9TritonScheduling.benchmark_fused_nodes.<locals>.<genexpr>G  s   è ø€ Ò:WÈA¸1¿:¹:¿<Ñ:Wùrú  )r]  )Úgenerate_kernel_code_from_nodesr$   r9  Úbenchmark_codegened_moduler   )r¿   rA  Ún_spills_thresholdrn  r]  s        rz   Úbenchmark_fused_nodesz&TritonScheduling.benchmark_fused_nodes?  sV   € ð
 ×7Ñ7¸ÐPTÐ7ÓUˆÜ×Ñ˜xÓ(ˆØ×.Ñ.ØÐ#´
Ñ:WÐQVÔ:WÓ0Wð /ó 
ð 	
r|   c                óâ  ‡‡	‡
‡‡‡— t        t        j                  j                  «      }t	        «       5  |j                  t        j                  j                  «       «      5  dŠˆfd„Š
ˆ
ˆfd„}ˆ
fd„}|�|nt        dg«      }t        j                  d|‰j                  «        |«       Š‰� ‰‰j                  fcddd«       cddd«       S ‰j                  «       Š	‰j                  Š‰j                  Š	  ‰ ‰j                  ‰	Ž d   «       ‰j(                  }t+        |«      d
k(  sJ ‚|d   j,                  |kD  rt'        d	«      ŠnNt/        j0                  ˆ	ˆˆfd„«      Št+        ‰j2                  «      dkD  r‰t/        j0                  ˆ	ˆfd„«      z
  Št        j                  d|‰«        |«        ‰‰j                  fcddd«       cddd«       S # t        $ rn}t         j"                  j$                  r‚ t        j                  d||«       t'        d	«      Š |«        ‰‰j                  fcY d}~cddd«       cddd«       S d}~ww xY w# 1 sw Y   nxY wddd«       y# 1 sw Y   yxY w)z$Benchmark an already compiled moduleNc                 ó~   •— ‰ j                   €J ‚t        j                  j                  ‰ j                   «      d   dz   S ©Nr   z.kernel_perf©Ú__file__ÚosÚpathÚsplitext©r]  s   €rz   Úcache_file_pathzDTritonScheduling.benchmark_codegened_module.<locals>.cache_file_pathU  s6   ø€ Ø—|‘|Ð/Ð/Ð/Ü—w‘w×'Ñ'¨¯©Ó5°aÑ8¸>ÑIÐIr|   c                 óŠ   •—  ‰«       } t        | d«      5 }|j                  t        ‰«      «       d d d «       y # 1 sw Y   y xY w)NÚw©ÚopenÚwriter´   )r…  Úfdrˆ  Úmss     €€rz   Ústore_cachez@TritonScheduling.benchmark_codegened_module.<locals>.store_cacheY  s;   ø€ Ù&Ó(�Ü˜$ “_ð &¨Ø—H‘HœS ›WÔ%÷&÷ &ñ &ús	   •9¹Ac                 óÆ   •—  ‰«       } t         j                  j                  | «      r.t        | «      5 }t	        |j                  «       «      cd d d «       S y # 1 sw Y   y xY wrq   )r„  r…  ÚexistsrŒ  ÚfloatÚread©r…  rŽ  rˆ  s     €rz   Ú
load_cachez?TritonScheduling.benchmark_codegened_module.<locals>.load_cache^  sM   ø€ Ù&Ó(�Ü—7‘7—>‘> $Ô'Ü˜d›ð 0 rÜ$ R§W¡W£YÓ/÷0ñ 0à÷0àús   ³AÁA Úunknownú%kernel src code for %s written to: %sr   z*Exception (%s) in compiling fused nodes %sÚinfrE   c                 ó4   •—  ‰ ‰j                   ‰ Ž d   «      S r=  ©Ú
clone_args©rC  ÚcallÚwrapped_jit_functions   €€€rz   r‹  z=TritonScheduling.benchmark_codegened_module.<locals>.<lambda>Ž  ó    ø€ ™DÐ!@Ð!5×!@Ñ!@À$Ð!GÈÑ!JÓK€ r|   c                 ó"   •—  ‰j                   ‰ Ž S rq   r›  ©rC  rŸ  s   €€rz   r‹  z=TritonScheduling.benchmark_codegened_module.<locals>.<lambda>–  s   ø€ Ð ?Ð 4× ?Ñ ?ÀÐ F€ r|   z+The fused kernel for %s took %.3f ms to run)r   rC   r  rç  r   r‡  rˆ  r   rþ  Údebugrƒ  Úget_argsrž  r`  rœ  Ú	Exceptionr   rŒ   Ú.disallow_failing_autotune_kernels_TESTING_ONLYr“  Ú	launchersr  Ún_spillsr'   Úbenchmark_gpurÉ  )r¿   r]  r}  r]  Údevice_interfacer�  r–  r  r§  rC  rˆ  rž  r�  rŸ  s    `       @@@@@rz   r|  z+TritonScheduling.benchmark_codegened_moduleJ  sK  ý€ ô 4´A·G±G×4GÑ4GÓHÐäÓ ñP	$à×#Ñ#¤A§G¡G×$GÑ$GÓ$IÓJñP	$ð ˆBôJõ&ô
ð )Ð4‘
¼*ÀiÀ[Ó:Qð ô �I‰IØ7ØØ—‘ôñ
 “ˆBØˆ~Ø˜3Ÿ<™<Ð'÷AP	$÷ P	$ò P	$ðD —<‘<“>ˆDØ—8‘8ˆDØ#&§;¡;Ð ð(ÙÐ4Ð)×4Ñ4°dÐ;¸AÑ>Ô?ð -×6Ñ6ˆIÜ�y“> QÒ&Ð&Ð&ð ˜‰|×$Ñ$Ð'9Ò9ä˜5“\‘ô !×.Ñ.ÝKó�ô Ð+×=Ñ=Ó>ÀÒBØœk×7Ñ7ÜFóñ �Bô �I‰IØ=ØØôñ
 ŒMØ�s—|‘|Ð#÷aP	$÷ P	$ò P	$øôP ò 
(Ü—=‘=×OÒOØÜ—	‘	Ø@ØØôô
 ˜5“\�Ù”Ø˜3Ÿ<™<Ð'Ô'÷eP	$÷ P	$ò P	$ûðP
(ú÷QP	$ð P	$ú÷ P	$÷ P	$ñ P	$úsh   ´.I%Á"AIÂ<	I%Ã(IÃ8GÄB3IÇ	I%Ç	IÇAIÈ0IÈ1IÈ5	I%ÉIÉIÉI	ÉI%É%I.c                óÄ  — |j                  d«      }|xr  t        d„ |j                  «       D «       «      }| j                  }|rddlm} |}|rd|d<   |j                  d«      r
d|d	<   d|d<   t        j                  |j                  «      s|j                  d	«      rJ ‚d|d	<   t        j                  j                  ||||«      } ||i |¤Ž}| j                  |||«      S )
NrS  c              3  ó<   K  — | ]  }|j                  «       –— Œ y ­wrq   )Úis_split_scan)r±   rC  s     rz   rÍ   z9TritonScheduling.create_kernel_choices.<locals>.<genexpr>¨  s   è ø€ ò (
Ø%)ˆD×Ñ× ñ(
ùrú  rE   )ÚTritonSplitScanKernelFÚoverride_cooperative_reductionr`  TÚoverride_persistent_reduction)Úcontains_oprÐ   Úscheduler_nodesrH  Útriton_split_scanr®  r”  rû  Úreduction_numelr^  rC   r±  Útriton_kernel_kwargsÚadd_multi_kernel_choices)	r¿   Úkernel_featuresÚkernel_argsÚkernel_kwargsÚis_scanr­  rH  r®  r  s	            rz   Úcreate_kernel_choicesz&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à$×(Ñ(Ð)HÔIÐIÐIØ=BˆMÐ9Ñ:äŸ	™	×6Ñ6Ø˜¨+°}ó
ˆñ ˜kÐ;¨]Ñ;ˆØ×,Ñ,¨V°[À-ÓPÐPr|   c           	     óÄ  — |g}t         j                  j                  s|S |j                  xr |j	                  d«       }|j
                  xr |j	                  d«       }|r%|j                   | j                  |i |¤ddi¤Ž«       |r›|j                  j                  }t        j                  j                  j                  |d«      r[|j                   | j                  |i |¤ddi¤Žx}«       |r2|j                  r&|j                   | j                  |i |¤dddœ¤Ž«       t        |«      dkD  r.|dd  D ]  }	|j                  |	_        Œ |j!                  d„ ¬«       |S )	Nr°  r¯  Fi   )r¯  r°  rE   c                ó   — | j                   S rq   )rÏ  )Úks    rz   r‹  z;TritonScheduling.add_multi_kernel_choices.<locals>.<lambda>ø  s   €  q×'=Ñ'=€ r|   )r2  )r   rŒ   Úmulti_kernelrÏ  r^  r©  rc  rH  r&  r´  rC   r  rù   rø  r  Úmust_keep_buffersr`  )
r¿   r  r¸  r¹  ÚkernelsÚoptional_persistentÚoptional_cooperativer[  rœ  Úkernel2s
             rz   r¶  z)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Ø—‘Ø-˜T×-Ñ-Ø$ðà'ñð 8=òð �Eôñ '¨5×+EÒ+EØ—N‘NØ(˜×(Ñ(Ø(ðà+ðð <AØ:?ó	ôô ˆw‹<˜!ÒØ" 1 2˜;ò E�à,2×,DÑ,D�Õ)ðEð �L‰LÑ=ˆLÔ>Øˆr|   c                óÒ  ‡‡‡‡‡‡‡— ˆfd„Šˆfd„}ˆˆˆfd„}dg }}d}t         j                  j                  }t        |«      t         j                  _        t         j                  j                  }t        |«      t         j                  _        t
        j                  dkD  }	t
        j                  dkD  }
| j                  |d|	|
d¬«      }|D �]Þ  \  }}}|D �cg c]  }|j                  «       ‘Œ }}|D ��cg c]  }|D ]  }|j                  «       ‘Œ Œ }}}|j                  t        t        j                  «      d«      }t        j                   |«      Št"        j%                  d	|‰j&                  «        |«       \  ŠŠ‰�&|‰z  }|‰z  }|j)                  ‰j&                  «       ŒÜ‰j+                  «       Š‰j,                  Š‰j.                  Š ‰ ‰j0                  ‰Ž d   «       ‰j2                  }t5        |«      d
k(  sJ ‚|d   j6                  dkD  rt9        d«      xŠŠn3t;        j<                  ˆˆˆfd„«      Št;        j<                  ˆˆfd„«      Št"        j%                  dt        d„ |D «       «      ‰‰«        |«        |‰z  }|‰z  }|j)                  ‰j&                  «       �Œá |t         j                  _        |t         j                  _        |||fS c c}w c c}}w )Nc                 ó~   •— ‰ j                   €J ‚t        j                  j                  ‰ j                   «      d   dz   S r�  r‚  r‡  s   €rz   rˆ  z@TritonScheduling.benchmark_combo_kernel.<locals>.cache_file_path   s6   ø€ Ø—<‘<Ð+Ð+Ð+Ü—7‘7×#Ñ# C§L¡LÓ1°!Ñ4°~ÑEÐEr|   c                 óð   •—  ‰«       } t         j                  j                  | «      rCt        | «      5 }t	        d„ |j                  «       j                  «       D «       «      cd d d «       S y# 1 sw Y   yxY w)Nc              3  ó2   K  — | ]  }t        |«      –— Œ y ­wrq   )r“  )r±   r  s     rz   rÍ   zNTritonScheduling.benchmark_combo_kernel.<locals>.load_cache.<locals>.<genexpr>  s   è ø€ Ò E¨a¤ q§Ñ Eùr  rB  )r„  r…  r’  rŒ  r¶  r”  r@  r•  s     €rz   r–  z;TritonScheduling.benchmark_combo_kernel.<locals>.load_cache  s^   ø€ Ù"Ó$ˆDÜ�w‰w�~‰~˜dÔ#Ü˜$“Zð F 2Ü Ñ E°2·7±7³9·?±?Ó3DÔ EÓE÷Fñ Fà÷Fàús   ³.A,Á,A5c                 ó¨   •—  ‰«       } t        | d«      5 }|j                  t        ‰«      dz   t        ‰«      z   «       d d d «       y # 1 sw Y   y xY w)NrŠ  r³  r‹  )r…  rŽ  rˆ  r�  Úms_clones     €€€rz   r�  z<TritonScheduling.benchmark_combo_kernel.<locals>.store_cache  sH   ø€ Ù"Ó$ˆDÜ�d˜C“ð 8 BØ—‘œ˜R› 3™¬¨X«Ñ6Ô7÷8÷ 8ñ 8ús   •*AÁAr   g        T)Úsubkernel_nodesÚcustom_part_algorithmÚenable_autotuneÚmixed_sizesÚonly_gen_src_coder`  r˜  rE   r™  c                 ó4   •—  ‰ ‰j                   ‰ Ž d   «      S r=  r›  r�  s   €€€rz   r‹  z9TritonScheduling.benchmark_combo_kernel.<locals>.<lambda>C  r   r|   c                 ó(   •—  ‰j                   ‰ Ž d   S r=  r›  r¢  s   €€rz   r‹  z9TritonScheduling.benchmark_combo_kernel.<locals>.<lambda>F  s   ø€ Ð;Ð0×;Ñ;¸TÐBÀ1ÑE€ r|   zDThe fused kernel for %s took %.3f ms to run, %.3f ms to clone inputsc              3  ó<   K  — | ]  }|j                  «       –— Œ y ­wrq   ry  rz  s     rz   rÍ   z:TritonScheduling.benchmark_combo_kernel.<locals>.<genexpr>K  s   è ø€ Ò<¨A˜1Ÿ:™:Ÿ<Ñ<ùrú  )rC   r  rÖ  r   Úinplaced_to_remover   Úcombo_kernels_autotuneÚcombo_kernel_allow_mixed_sizesÚgenerate_combo_kernel_codeÚ	get_nodesrX  rh  r´   r8   rŽ  r$   r9  rþ  r£  rƒ  rc  r¤  rž  r`  rœ  r§  r  r¨  r“  r'   r©  )r¿   Ú	node_listr–  r�  Útotal_msÚ	file_listÚtotal_clone_msÚremoved_buffers_origÚinplaced_to_remove_origrÍ  rÎ  Úkernel_code_listrn  rŒ  Ú
node_grouprC  Úfused_node_listsrA  r  Únamesr§  rC  rˆ  rž  r]  r�  rÊ  rŸ  s                        @@@@@@@rz   Úbenchmark_combo_kernelz'TritonScheduling.benchmark_combo_kernelû  s«  þ€ ô
	Fô	 ö	8ð
   �)ˆØ #ˆÜ Ÿw™w×6Ñ6ÐÜ",Ð-AÓ"BŒ�‰ÔÜ"#§'¡'×"<Ñ"<ÐÜ%/Ð0GÓ%HŒ�‰Ô"Ü ×7Ñ7¸!Ñ;ˆÜ×;Ñ;¸aÑ?ˆØ×:Ñ:Ø%Ø"&Ø+Ø#Ø"ð ;ó 
Ðð (8ó 2	+Ñ#ˆH�a˜Ø=GÖH°T §¡Õ 0ÐHÐÐHØ/?×O eÈÒOÀA�Q—Z‘Z•\ÐO�\ÐOˆEÑOà×'Ñ'¬¬K×,CÑ,CÓ(DÀiÓPˆHÜ×"Ñ" 8Ó,ˆCä�I‰IØ7ØØ—‘ôñ
 &›<‰LˆB�Øˆ~Ø˜B‘�Ø (Ñ*�Ø× Ñ  §¡Ô.Øà—<‘<“>ˆDØ—8‘8ˆDØ#&§;¡;Ð ñ Ð0Ð%×0Ñ0°$Ð7¸Ñ:Ô;à,×6Ñ6ˆIÜ�y“> QÒ&Ð&Ð&Ø˜‰|×$Ñ$ qÒ(ä % e£Ð,�‘Xô !×.Ñ.ÝKó�ô '×4Ñ4ÜEó�ô �I‰IØVÜÑ<°Ô<Ó<ØØô	ñ ŒMØ˜‰NˆHØ˜hÑ&ˆNØ×Ñ˜SŸ\™\Ö*ðe2	+ðf #7Œ�‰ÔØ%<Œ�‰Ô"Ø˜¨Ð2Ð2ùòi  IùÛOs   ÃKÃ7K#)rL  zOptional[Scheduler]r~   r   )r‡  ztorch.device)rì  )r~   útuple[float, str])rì  N)r]  zOptional[OrderedSet[str]]r~   rã  )r·  rg   r¸  ú	list[Any]r¹  r  r~   úlist[TritonKernel])r  r”  r¸  rä  r¹  r  r~   rå  )rr   r€   r�   r”  rH  rƒ   r   rH   ÚFOREACHÚ	BUCKETIZEÚINPLACE_BUFFERSÚMASKED_SCATTER_WITH_INDEXÚSCANÚSORTÚTRITON_TEMPLATESÚTUPLE_REDUCTIONrP  r  r„   rR  r^  rk  r~  r|  r»  r¶  râ  r  r  s   @rz   rG  rG  É  s  ø… Ø)€K�Ó)Ù!à×"Ñ"Ø×$Ñ$Ø×*Ñ*Ø×4Ñ4Ø×ÑØ×ÑØ×+Ñ+Ø×*Ñ*ð		
óÐõ:ð ò$ó ð$òò48ôt	
ð RVðU$Ø5NðU$à	óU$ðn#Qà+ð#Qð ð#Qð &ð	#Qð
 
ó#QðJ3àð3ð ð3ð &ð	3ð
 
ó3öjZ3r|   rG  c                ó  — g }| j                  «       }|�t        |t        j                  «      sJ ‚|r0|j                  €$|j                  | j                  «       › d�«       |S ddlm} | j                  «       }|€J ‚| j                  j                  |«      }t        |t        |f«      sJ dt        |«      › �«       ‚t        j                  j!                  |«      5  t"        j$                  }|j'                  | j)                  «       «      j+                  «       }|t"        _        d d d «       |j                  | j                  «       › d�«       |j                  t-        j.                  d«      «       |S # 1 sw Y   ŒRxY w)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_noderb  r   ÚMultiTemplateBufferÚmake_kernel_renderrc  rX  Ú0torch._inductor.codegen.cuda_combined_schedulingrï  r…  rL  Úget_backendr[   rµ  rC   r  Úset_current_devicer    Úgenerated_kernel_countr{  r×  r  rž  rf  )rC  ÚlinesÚmulti_templaterï  r‡  ÚbackendÚold_generated_kernel_countÚtriton_codes           rz   rJ  rJ  X  sd  € Ø€EØ×+Ñ+Ó-€NØÐ!¤Z°Ä×@VÑ@VÔ%WÐWÐWÙ˜.×;Ñ;ÐCØ�‰˜Ÿ™›Ð(Ð(JÐKÔLð2 €Lõ/	
ð —‘Ó"ˆØÐ!Ð!Ð!Ø—.‘.×,Ñ,¨VÓ4ˆÜ˜'¤NÐ4JÐ#KÔLð 	
ØkÔlpÐqxÓlyÐkzÐ{ó	
ÐLô �W‰W×'Ñ'¨Ó/ñ 	Hô *1×)GÑ)GÐ&Ø!×AÑAØ—‘Ó óç‰e‹gð ð .HŒGÔ*÷	Hð 	�‰˜Ÿ™›Ð(¨Ð6Ô7Ø�‰”X—_‘_ [°&Ó9Ô:Ø€L÷	Hð 	Hús   ÃA	E6Å6E?rÜ   )r  r´   rd  rå   re  rå   r~   r´   rA  )rè  r  r~   r6  )rè  r  r~   r  )rè  r  r~   r}   )rý  zUnion[CSEVariable, Any]r~   r}   )r~   re   r[  )rx   r}   r~   zCallable[[_T], _T])rC  r.   r~   rW  )¼Ú
__future__r   rs  r\  r  rä  r  Úloggingr  r„  rž  Úcollections.abcr   r   r   Útypingr   r   r	   r
   r   r   r®   Úsympy.printing.precedencer   rs   Ú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   Úutils._sympy.symbolr   r   r   r   Úutils._sympy.value_rangesr   r‡   r   r   r    ri  r!   Ú	codecacher"   r#   r$   Úops_handlerr%   Úruntimer&   Úruntime.benchmarkingr'   Úruntime.hintsr(   r)   r*   r+   Úruntime.runtime_utilsr,   r-   rL  r.   r/   r0   r1   r2   r3   r4   r5   r6   r7   r8   r9   r:   r;   r<   r=   r>   r?   Úvirtualizedr@   rx  rA   rB   rC   Úwrapper_benchmarkrD   Úblock_analysisrF   ÚcommonrG   rH   rI   rJ   rK   rL   rM   rN   rO   rP   rQ   rR   rS   rT   rU   ÚsimdrV   rW   rX   rY   rZ   r[   Útriton_utilsr\   r]   r^   r_   r`   r  ra   Útypesrb   rc   r  re   rf   Úsimd_kernel_featuresrg   rh   Ú	getLoggerrr   rþ  Ú_loggingÚgetArtifactLoggerÚperf_hint_logÚschedule_logÚ
fusion_logrm   r�   r”   r–   Ú	dataclassr³   rß   r   rl  r§  r2  ré  rî  rñ  rö  rø  rû  rþ  r   r  r+  r-  Ú_initialize_pointwise_overridesr   rU  rà   ro  rˆ  r´   r¶  r�  r”  rG  rJ  r…   r|   rz   ú<module>r"     s:  ðå "ã Û Û Û Û Û Û Û 	Û ß .Ý ß F× Fã Ý 0ã Û ß $Ð $Ý Cß <Ý 0Ý /ß KÑ KÝ 2ç XÓ XÝ 4ß "Ñ "Ý (ß 8Ñ 8Ý (Ý 'Ý .÷ó ÷ Dß WÓ W÷÷ ÷ ÷ ÷  CÓ BÝ BÝ /÷÷ ÷ ÷ ñ ÷"÷ ÷õ õ %ñ Ý ÝåLåÝ8á	�‹€Bà€g×Ò˜Ó!€Ø—’×0Ò0°¸<ÓH€Ø�~Š~×/Ò/°¸*ÓE€Ø�^Š^×-Ò-¨h¸ÓA€
Ù“€÷6ñ 6ñ  ˆ4ƒòó ðñ$ ˆ4ƒòó ð÷*,ñ ,ð: ×Ò÷Rð Ró ðRð6 ×Ò÷N+ð N+ó ðN+ðb+Øð+Ø/ð+Ø<Pð+àó+ô>aQ�Mô aQñH 	‹×Ò€ó3ó
óó&ó8ó
;óPô˜ô ó*(ô.ôbs&�kô s&ðl × /Ò /°Ô 9ôD$˜Oô D$÷N$+ñ $+ðN ×Ò÷:ð :ó ð:÷&!ñ !ðH ×Ò÷#ð #ó ð#ô
�Ð% u¨S°%¸¸S¸±/Ð-AÑ'BÐBÑCô 
ôU&A�:Ð/Ñ0ô U&AôpLL3�~ô L3ô^r|   