Ë
    1²jx  ã                   óþ   — d dl mZmZmZmZmZ d dlZd dlmZ ddl	m
Z ddlmZmZmZmZmZmZmZ d dlmZmZ d dlmZmZ  ed	«      Zefd
edeg ef   fd„Zd„ Zdd„Z de _!         G d„ d«      Z" G d„ dee   «      Z#y)é    )ÚSequenceÚListÚTypeVarÚTupleÚCallableN)ÚTritonSemanticé   )Ú_core)Ú
AutoLayoutÚDistributedLayoutÚDistributedLinearLayoutÚSliceLayoutÚSharedLayoutÚCoalescedLayoutÚSharedLinearLayout)ÚGluonOpBuilderÚcompute_tmem_reg_layout)Úflatten_values_to_irÚunflatten_ir_valuesÚTensorTyÚcondÚmsg_fnc                 ó"   — | s | |«       «      ‚y ©N© )r   r   Úcategorys      úk/var/www/html/ia-enhanced/venv/lib/python3.12/site-packages/triton/experimental/gluon/language/_semantic.pyÚ_checkr      s   € ÙÙ‘v“xÓ Ð ð ó    c                 óJ   — t        | t        «      xr t        d„ | D «       «      S )Nc              3   ó<   K  — | ]  }t        |t        «      –— Œ y ­wr   ©Ú
isinstanceÚint)Ú.0Úis     r   ú	<genexpr>z_is_int_list.<locals>.<genexpr>   s   è ø€ Ò.QÀa¬z¸!¼S×/AÑ.Qùó   ‚)r#   r   Úall©Úvalues    r   Ú_is_int_listr,      s   € Ü�eœXÓ&ÒQ¬3Ñ.QÈ5Ô.QÓ+QÐQr   c                 ó@  ‡‡‡‡‡‡— t        t        ‰t        «      d„ «       t        ‰dv ˆfd„«       t        t        ‰t        «      ˆfd„«       t        ‰dk\  xr ‰‰dz
  z  dk(  d„ «       t	        ‰«      Št        t        d	„ ‰D «       «      ˆfd
„«       t        ‰«      }t        |dk(  d„ «       |€g }‰dk(  }|rdn‰Š|r |D ]  }t        t        |«      |k(  d„ «       Œ t        | ‰|‰‰|«      Št        ‰d uˆˆˆfd„«       |�r‰d   }	‰j                  sWt        ‰j                  d   d|	dz  gk(  ˆfd„«       ‰j                  j                  d|	dz  g«       ddg‰j                  d<   ‰S ‰j                  d   d|	dz  gk7  r–| j                  Šdt        ‰j                  «      z  }
t        |
d‰z  kD  ˆfd„«       ‰j                  }dD ]B  }t        ‰|«      }t        |«      D ]&  \  }}|d|	dz  gk(  sŒ||   |d   c|d<   ||<   ‰c c S  ŒD J d|› �«       ‚‰S )Nc                   ó   — y)Nzinstr_variant must be a stringr   r   r   r   ú<lambda>z*_compute_tmem_reg_layout.<locals>.<lambda>   ó   � r   )Ú32x32bÚ16x64bÚ16x128bÚ16x256bÚ16x32bx2Ú32x32b_splitnc                  ó   •— d‰ › �S )Nzunknown instr_variant: r   )Úinstr_variants   €r   r/   z*_compute_tmem_reg_layout.<locals>.<lambda>   s   ø€ Ð,¨]¨OÐ<€ r   c                  ó    •— dt        ‰ «      ›�S )Nz!num_warps must be an int but got ©Útype©Ú	num_warpss   €r   r/   z*_compute_tmem_reg_layout.<locals>.<lambda>   s   ø€ Ð1RÔSWÐXaÓSbÐReÐ/f€ r   é   r	   r   c                   ó   — y)Nz)num_warps must be a power of two and >= 4r   r   r   r   r/   z*_compute_tmem_reg_layout.<locals>.<lambda>   r0   r   c              3   ó<   K  — | ]  }t        |t        «      –— Œ y ­wr   r"   )r%   Údims     r   r'   z+_compute_tmem_reg_layout.<locals>.<genexpr>   s   è ø€ Ò5¨Œz˜#œs×#Ñ5ùr(   c                  ó   •— d‰ › �S )Nz#shape entries must be ints but got r   ©Úshapes   €r   r/   z*_compute_tmem_reg_layout.<locals>.<lambda>   s   ø€ ÐAdÐejÐdkÐ?l€ r   é   c                   ó   — y)Nzexpected a 2D tensorr   r   r   r   r/   z*_compute_tmem_reg_layout.<locals>.<lambda>   r0   r   r6   r1   c                   ó   — y)Nzcga_layout basis rank mismatchr   r   r   r   r/   z*_compute_tmem_reg_layout.<locals>.<lambda>(   r0   r   c                  ó   •— d‰ › d‰› d‰› �S )NzTMEM layout 'z' unsupported for shape z and num_warps r   )Úatom_variantr=   rD   s   €€€r   r/   z*_compute_tmem_reg_layout.<locals>.<lambda>3   s   ø€ �] < .Ð0HÈÈÈÐ_hÐ^iÐj€ r   éÿÿÿÿc                  ó   •— d‰ › �S )NzJsplitn with 1 register requires the last lane basis to be [0, N / 2]. Got r   )Ú
layout_objs   €r   r/   z*_compute_tmem_reg_layout.<locals>.<lambda>;   s   ø€ ÐgÐhrÐgsÐt€ r   é    c                  ó   •— dd‰ z  › d‰ › d�S )NzETo be able to `tmem.load` into `tl.split` you need to have more than rM   ú z¥-bit registers, as you need to use the instruction 32x32b.x1 twice. You can always load into instr_variant="32x32b" and then convert_layout to this layout otherwise.r   )Úbitwidths   €r   r/   z*_compute_tmem_reg_layout.<locals>.<lambda>B   s&   ø€ ð 3Ø˜h™Ð' q¨¨
ð 3]ð3]€ r   )Ú
lane_basesÚ
warp_basesz6splitn requires at least one basis of [0, N / 2]. Got )r   r#   Ústrr$   Úlistr)   Úlenr   Ú	reg_basesrQ   ÚappendÚprimitive_bitwidthÚgetattrÚ	enumerate)Ú
element_tyrD   Úlayoutr=   r8   Ú
cga_layoutÚrankÚsplitnÚbasisÚNÚnum_regrV   Ú	bases_strÚbasesr&   rI   rP   rL   s    ` ``          @@@r   Ú_compute_tmem_reg_layoutre      sj  ý€ Ü
Œ:�m¤SÓ)Ñ+SÔTÜ
ˆ=ÐcÐcÛ<ô>ä
Œ:�i¤Ó%Ó'fÔgÜ
ˆ9˜‰>Ò@˜y¨I¸©MÑ:¸qÑ@ÑBuÔvä�‹K€EÜ
Œ3Ñ5¨uÔ5Ó5Ó7lÔmÜˆu‹:€DÜ
ˆ4�1‰9Ñ4Ô5àÐØˆ
Ø˜oÑ-€FÙ%‘8¨=€LáØò 	QˆEÜ”3�u“: Ñ%Ñ'OÕPð	Qô )ØØØØØØó€Jô ˆ:˜TÐ!Ýjôlò Ø�!‰HˆØ×#Ò#ô �:×(Ñ(¨Ñ,°°A¸±F°Ñ;Ûtôvà× Ñ ×'Ñ'¨¨A°©F¨Ô4Ø)*¨A¨ˆJ×!Ñ! "Ñ%ð$ Ðð# ×!Ñ! "Ñ%¨!¨Q°!©V¨Ò4Ø!×4Ñ4ˆHØœ˜Z×1Ñ1Ó2Ñ2ˆGÜØ˜" ™.Ñ(ó +]ô^ð #×,Ñ,ˆIØ9ò *�	Ü 
¨IÓ6�Ü )¨%Ó 0ò *‘H�A�uØ  A¨¡F Ó+Ø27¸±(¸IÀb¹MÐ/˜	 "™ u¨Q¡xØ)Ô)ñ*ð*ð \ÐRÐSYÐRZÐ[Ó[�5ØÐr   Tc                   ó$   — e Zd Zdefd„Zd„ Zd„ Zy)ÚGluonCallerContextr=   c                 ó   — || _         y r   r<   )Úselfr=   s     r   Ú__init__zGluonCallerContext.__init__W   s	   € Ø"ˆ�r   c                 ó    — d| j                   › �S )NÚ_NWr<   ©ri   s    r   ÚmanglezGluonCallerContext.mangleZ   s   € Ø�T—^‘^Ð$Ð%Ð%r   c                 óZ   — |j                  d|j                  | j                  «      «       y )Nzttg.num-warps)Úset_attrÚget_int32_attrr=   )ri   ÚfnÚbuilders      r   Úinitialize_calleez$GluonCallerContext.initialize_callee]   s   € Ø
�‰�O W×%;Ñ%;¸D¿N¹NÓ%KÕLr   N)Ú__name__Ú
__module__Ú__qualname__r$   rj   rn   rt   r   r   r   rg   rg   U   s   „ ð# #ó #ò&óMr   rg   c            
       ó|  ‡ — e Zd ZU ej                  ZeZeed<   defd„Zd„ Z	d„ Z
dee   dee   fd„Zded	ed
efd„Zdeded
efˆ fd„Zded
eeef   fˆ fd„Zdedee   d
efˆ fd„Zdedee   d
efd„Zdeded
efˆ fd„Zˆ fd„Zdedee   defˆ fd„Zd„ Zd„ Zd@d„Zd„ Zd„ Zd „ Zd!„ Zd"„ Z d#„ Z!d$„ Z"d%„ Z#d&„ Z$d'„ Z%d(„ Z&d)„ Z'd*„ Z(d+„ Z)d,„ Z*e+d-„ «       Z,d.e-e   d	ed/ed
eed0f   fd1„Z.d.e-e   d	ed
eed0f   fd2„Z/ded3ed4ed
efd5„Z0dededed
efd6„Z1d7ed8ed	ed
efd9„Z2d7ed
efd:„Z3d;e-e   d<e-e   fd=„Z4d>„ Z5d?„ Z6ˆ xZ7S )AÚGluonSemanticrs   c                 ó   — || _         y r   )rs   )ri   rs   s     r   rj   zGluonSemantic.__init__g   s	   € Øˆ�r   c                 ó–   — |g k(  r|}n0t        j                  ||| j                  j                  |«      «      }| j	                  ||«      S r   )ÚttglÚdistributed_typers   Úget_gluon_layout_from_tensorÚtensor)ri   ÚhandleÚ	scalar_tyrD   Útys        r   Ú_wrap_handle_infer_layoutz'GluonSemantic._wrap_handle_infer_layoutj   sC   € Ø�BŠ;Ø‰Bä×&Ñ& y°%¸¿¹×9bÑ9bÐciÓ9jÓkˆBØ�{‰{˜6 2Ó&Ð&r   c                 óx   — | j                  |j                  |j                  j                  |j                  «      S r   )rƒ   r€   r;   ÚscalarrD   )ri   r   s     r   Ú_wrap_tensor_infer_layoutz'GluonSemantic._wrap_tensor_infer_layoutq   s,   € Ø×-Ñ-¨f¯m©m¸V¿[¹[×=OÑ=OÐQW×Q]ÑQ]Ó^Ð^r   Ú	lhs_shapeÚ	rhs_shapec                 óX  — t        |«      t        |«      k7  rt        d|› d|› �«      ‚g }t        |«      D ]q  \  }}||   }|dk(  r|j                  |«       Œ"|dk(  s||k(  r|j                  |«       Œ>t        dt	        |«      z   dz   t	        |«      z   dz   t	        |«      z   «      ‚ |S )Nú!Cannot broadcast, rank mismatch: ú, r	   z?Cannot make_shape_compatible: incompatible dimensions at index ú: ú and )rU   Ú
ValueErrorrZ   rW   rS   )ri   r‡   rˆ   Ú	ret_shaper&   ÚleftÚrights          r   Ú_broadcast_shapeszGluonSemantic._broadcast_shapest   sØ   € Üˆy‹>œS ›^Ò+ÜÐ@ÀÀÈ2ÈiÈ[ÐYÓZÐZàˆ	Ü  Ó+ò 	a‰GˆAˆtØ˜a‘LˆEØ�qŠyØ× Ñ  Õ'Ø˜1’* %¨4¢-Ø× Ñ  Õ&ä ð "-Ü/2°1«vñ"6Ø8<ñ"=Ü?BÀ4»yñ"IØKRñ"SÜUXÐY^ÓU_ñ"`ó að að	að Ðr   ÚinputÚaxisÚreturnc                 óª  ‡‡‡— ‰j                   D �cg c]  }t        j                  |«      ‘Œ }}|j                  ‰d«       ‰dk  r‰t	        ‰j                   «      z  Št        t        ‰j                  t        j                  «      ˆfd„«       ‰j                  j                  Št        t        ‰t        t        t        f«      ˆfd„«       t        t        ‰t        t        f«      xs ‰j                  ‰k(  ˆˆfd„«       | j                  j                  ‰j                   ‰«      }| j#                  |‰j                  j$                  |«      S c c}w )Nr	   r   c                  ó"   •— d‰ j                   ›�S ©Nz=expected expand_dims input to be a distributed_type but got: r:   ©r“   s   €r   r/   z+GluonSemantic.expand_dims.<locals>.<lambda>Œ   ó   ø€ ÐVÐW\×WaÑWaÐVdÐe€ r   c                  ó   •— d‰ › �S )Nz;expected expand_dims input to have a SliceLayout, but got: r   ©r\   s   €r   r/   z+GluonSemantic.expand_dims.<locals>.<lambda>�   s   ø€ ÐTÐU[ÐT\Ð]€ r   c                  ó(   •— d‰ › d‰j                   › �S )Nz7expected expand_dims input layout to be sliced in axis z	 but got ©rA   )r”   r\   s   €€r   r/   z+GluonSemantic.expand_dims.<locals>.<lambda>’   s   ø€ ÐMÈdÈVÐS\Ð]c×]gÑ]gÐ\hÐi€ r   )rD   r|   Ú_unwrap_if_constexprÚinsertrU   r   r#   r;   r}   r\   r   r   r   rA   rs   Úcreate_expand_dimsr€   rƒ   r…   )ri   r“   r”   ÚxÚ	dst_shaper€   r\   s    ``   @r   Úexpand_dimszGluonSemantic.expand_dims„   s  ú€ Ø;@¿;¹;ÖG°a”T×.Ñ.¨qÕ1ÐGˆ	ÐGØ×Ñ˜˜qÔ!à�!Š8Ø”C˜Ÿ™Ó$Ñ$ˆDäŒz˜%Ÿ*™*¤d×&;Ñ&;Ó<Ûeô	gà—‘×"Ñ"ˆÜŒz˜&¤;´
¼OÐ"LÓMÛ]ô	_äÜ�v¤
¬OÐ<Ó=ÒSÀÇÁÈtÑASÜiô	kð —‘×0Ñ0°·±¸tÓDˆØ×-Ñ-¨f°e·j±j×6GÑ6GÈÓSÐSùò! Hs   ’EÚaÚbc                 ó¤   •— | j                  ||«      \  }}t        |j                  g k7  d„ «       t        ‰| �  ||«      }| j                  |«      S )Nc                   ó   — y)NzCannot join scalars in gluonr   r   r   r   r/   z$GluonSemantic.join.<locals>.<lambda>™   r0   r   )Úbroadcast_impl_valuer   rD   ÚsuperÚjoinr†   )ri   r¥   r¦   r+   Ú	__class__s       €r   r«   zGluonSemantic.join—   sM   ø€ Ø×(Ñ(¨¨AÓ.‰ˆˆ1Üˆq�w‰w˜"‰}ÑDÔEÜ‘‘˜Q Ó"ˆØ×-Ñ-¨eÓ4Ð4r   c                 ól   •— t         ‰| �  |«      \  }}| j                  |«      | j                  |«      fS r   )rª   Úsplitr†   )ri   r¥   ÚlhsÚrhsr¬   s       €r   r®   zGluonSemantic.split�   s7   ø€ Ü‘7‘= Ó#‰ˆˆSØ×-Ñ-¨cÓ2°D×4RÑ4RÐSVÓ4WÐWÐWr   Údimsc                 óF   •— t         ‰| �  ||«      }| j                  |«      S r   )rª   Úpermuter†   )ri   r“   r±   r+   r¬   s       €r   r³   zGluonSemantic.permute¡   s$   ø€ Ü‘‘  tÓ,ˆØ×-Ñ-¨eÓ4Ð4r   rD   c                 ó’  ‡‡‡— t        t        ‰j                  t        j                  «      ˆfd„«       ‰j                  j                  «       Št        t        ‰«      t        ‰«      k(  ˆˆfd„«       ‰‰k(  r‰S t        ‰«      D ]0  \  }}‰|   |k7  sŒ|dk7  sŒt        d‰|   › d|› d|› d‰› d‰› �
«      ‚ t        j                  ‰j                  j                  ‰‰j                  j                  «      }| j                  j                  ‰j                  |j                  | j                  «      «      }| j                  ||«      S )	Nc                  ó"   •— d‰ j                   ›�S r˜   r:   r™   s   €r   r/   z4GluonSemantic.broadcast_impl_shape.<locals>.<lambda>§   rš   r   c                  ó   •— d‰› d‰ › �S )NrŠ   r‹   r   )rD   Ú	src_shapes   €€r   r/   z4GluonSemantic.broadcast_impl_shape.<locals>.<lambda>©   s   ø€ Ð7XÐYbÐXcÐceÐfkÐelÐ5m€ r   r	   z3Cannot broadcast, the expanded size of the tensor (z ) must match the existing size (z) at non-singleton dimension rŒ   r‹   )r   r#   r;   r|   r}   Úget_block_shapesrU   rZ   rŽ   r…   r\   rs   Úcreate_broadcastr€   Úto_irr   )ri   r“   rD   r&   ÚitemÚret_tyr€   r·   s    ``    @r   Úbroadcast_impl_shapez"GluonSemantic.broadcast_impl_shape¥   s-  ú€ ÜŒz˜%Ÿ*™*¤d×&;Ñ&;Ó<Ûeô	gà—J‘J×/Ñ/Ó1ˆ	ÜŒs�9‹~¤ U£Ñ+Ô-mÔnØ�IÒØˆLÜ  Ó+ò 	@‰GˆAˆtØ�Q‰x˜4Ó D¨A£IÜ Ð#VÐW\Ð]^ÑW_ÐV`ð aCØCGÀ&ð I%Ø%& C r¨)¨°B°u°gð"?ó @ð @ð	@ô
 ×&Ñ& u§z¡z×'8Ñ'8¸%ÀÇÁ×ARÑARÓSˆØ—‘×.Ñ.¨u¯|©|¸V¿\¹\È$Ï,É,Ó=WÓXˆØ�{‰{˜6 6Ó*Ð*r   r¯   r°   c                 óP  •‡‡	— |j                   Š|j                   Š	‰j                  «       r‰	j                  «       st        ‰
| �  ||«      S t	        t        ‰t        j                  «      ˆfd„«       t	        t        ‰	t        j                  «      ˆ	fd„«       ‰j                  «       }‰	j                  «       }| j                  ||«      }t        ‰j                  t        «      }t        ‰	j                  t        «      }|r|s| j                  |‰	j                  «      }n_|r|s| j                  |‰j                  «      }n>‰j                  ‰	j                  k7  r%t        d‰j                  › d‰	j                  › �«      ‚| j                  ||«      }| j                  ||«      }||fS )Nc                  ó   •— d‰ ›�S )Nz@expected broadcast left input to be a distributed_type but got: r   )Úlhs_tys   €r   r/   z4GluonSemantic.broadcast_impl_value.<locals>.<lambda>½   s   ø€ ÐYÐZ`ÐYcÐd€ r   c                  ó   •— d‰ ›�S )NzAexpected broadcast right input to be a distributed_type but got: r   )Úrhs_tys   €r   r/   z4GluonSemantic.broadcast_impl_value.<locals>.<lambda>¿   s   ø€ ÐZÐ[aÐZdÐe€ r   zLayout mismatch in broadcast: z vs )r;   Úis_blockrª   r©   r   r#   r|   r}   r¸   r’   r\   r   Úset_auto_layoutrŽ   r½   )ri   r¯   r°   r‡   rˆ   r�   Úis_lhs_autoÚis_rhs_autorÀ   rÂ   r¬   s           @@€r   r©   z"GluonSemantic.broadcast_impl_valueµ   sX  ú€ Ø—‘ˆØ—‘ˆà�‰Ô ¨¯©Ô(9Ü‘7Ñ/°°SÓ9Ð9äŒz˜&¤$×"7Ñ"7Ó8Ûdô	fäŒz˜&¤$×"7Ñ"7Ó8Ûeô	gð ×+Ñ+Ó-ˆ	Ø×+Ñ+Ó-ˆ	Ø×*Ñ*¨9°iÓ@ˆ	ä  §¡´
Ó;ˆÜ  §¡´
Ó;ˆÙ™{Ø×&Ñ& s¨F¯M©MÓ:‰CÙ¡Ø×&Ñ& s¨F¯M©MÓ:‰CØ�]‰]˜fŸm™mÒ+ÜÐ=¸f¿m¹m¸_ÈDÐQW×Q^ÑQ^ÐP_Ð`ÓaÐaà×'Ñ'¨¨YÓ7ˆØ×'Ñ'¨¨YÓ7ˆØ�Cˆxˆr   c                 ó–   •— ||z
  g}|€
t        «       }t        j                  t        j                  ||«      }t        ‰| �  |||¬«      S )N)r¼   )r   r|   r}   Úint32rª   Úarange)ri   ÚstartÚendr\   rD   r¼   r¬   s         €r   rÉ   zGluonSemantic.arangeÒ   sG   ø€ Ø�u‘�ˆØˆ>Ü“\ˆFÜ×&Ñ&¤t§z¡z°5¸&ÓAˆÜ‰w‰~˜e S°ˆ~Ó8Ð8r   r£   Úcan_reorderc                 ód   •— t        | d„ «       t        ‰| �	  |||«      }| j                  |«      S )Nc                   ó   — y)Nz%can_reorder is not supported in gluonr   r   r   r   r/   z'GluonSemantic.reshape.<locals>.<lambda>Ú   r0   r   )r   rª   Úreshaper†   )ri   r“   r£   rÌ   r+   r¬   s        €r   rÏ   zGluonSemantic.reshapeÙ   s3   ø€ Ü�;ˆÑ OÔPÜ‘‘  y°+Ó>ˆØ×-Ñ-¨eÓ4Ð4r   c                 ó  — t        |«      dk(  r|S t        j                  |j                  ||«      }| j                  j                  |j                  | j                  «      |j                  «      }t        j                  ||«      S )Nr   )	rU   r|   r}   Údtypers   Úcreate_splatrº   r€   r   )ri   r+   rD   r\   r¼   r€   s         r   ÚsplatzGluonSemantic.splatÞ   sb   € Üˆu‹:˜Š?ØˆLÜ×&Ñ& u§{¡{°E¸6ÓBˆØ—‘×*Ñ*¨6¯<©<¸¿¹Ó+EÀuÇ|Á|ÓTˆÜ�{‰{˜6 6Ó*Ð*r   c                 ód   — | j                  ||«      }|€
t        «       }| j                  |||«      S r   )Úmake_scalarr   rÓ   )ri   rD   r+   rÑ   r\   r…   s         r   ÚfullzGluonSemantic.fullå   s3   € Ø×!Ñ! %¨Ó/ˆØˆ>Ü“\ˆFØ�z‰z˜& %¨Ó0Ð0r   c                 óÌ  ‡‡— |j                   Št        t        ‰t        j                  «      ˆfd„«       t        t        ‰t        j
                  «      ˆfd„«       t        j                  ‰j                  ‰j                  ‰«      }|j                  | j                  «      }|r…| j                  j                  ||j                  «      s_t        d‰j                  › d‰› d| j                  ‰j                  ‰j                  «      › d| j                  ‰‰j                  «      › �«      ‚| j                  j                  ||j                  «      }t        j                   ||«      S )Nc                  ó   •— d‰ ›�S )Nz@expected convert_layout input to be a distributed_type but got: r   )r‚   s   €r   r/   z.GluonSemantic.convert_layout.<locals>.<lambda>î   s   ø€ ÐYÐZ\ÐY_Ð`€ r   c                  ó   •— d‰ › �S ©Nz4expected 'layout' to be a DistributedLayout but got r   rœ   s   €r   r/   z.GluonSemantic.convert_layout.<locals>.<lambda>ð   ó   ø€ ÐMÈfÈXÐV€ r   zlayout conversion from z to z) is not trivial.
The linear layouts are:
ú
)r;   r   r#   r|   r}   r   r[   rD   rº   rs   Úis_convert_layout_trivialr€   Ú	TypeErrorr\   Úto_linear_layoutÚcreate_convert_layoutr   )ri   r+   r\   Úassert_trivialr¼   Ú	ret_ty_irr€   r‚   s     `    @r   Úconvert_layoutzGluonSemantic.convert_layoutë   s  ù€ Ø�Z‰ZˆÜŒz˜"œd×3Ñ3Ó4Û`ô	bäŒz˜&¤$×"8Ñ"8Ó9ÛVô	Xä×&Ñ& r§}¡}°b·h±hÀÓGˆØ—L‘L §¡Ó.ˆ	Ù $§,¡,×"HÑ"HÈÐTY×T`ÑT`Ô"aÜÐ5°b·i±i°[ÀÀVÀHð M8Ø8<×8MÑ8MÈbÏiÉiÐY[×YaÑYaÓ8bÐ7cÐceØ#×4Ñ4°V¸R¿X¹XÓFÐGðIó Jð Jð —‘×3Ñ3°I¸u¿|¹|ÓLˆÜ�{‰{˜6 6Ó*Ð*r   c                 ó"  ‡‡‡— t        t        ‰t        j                  «      ˆfd„«       t        t	        ‰«      ˆfd„«       t        t        ‰t        j
                  «      ˆfd„«       t        j                  ‰‰‰‰«      }|�@| j                  j                  |j                  | j                  «      |j                  «      }n4| j                  j                  |j                  | j                  «      «      }t        j                  |‰‰‰‰«      S )Nc                  ó   •— d‰ › �S )Nz,expected 'element_ty' to be a dtype but got r   )r[   s   €r   r/   z/GluonSemantic.allocate_shared.<locals>.<lambda>û   s   ø€ Ð=iÐjtÐiuÐ;v€ r   c                  ó   •— d‰ › �S ©Nz1all elements of 'shape' must be integers but got r   rC   s   €r   r/   z/GluonSemantic.allocate_shared.<locals>.<lambda>ü   ó   ø€ Ð._Ð`eÐ_fÐ,g€ r   c                  ó   •— d‰ › �S ©Nz/expected 'layout' to be a SharedLayout but got r   rœ   s   €r   r/   z/GluonSemantic.allocate_shared.<locals>.<lambda>þ   ó   ø€ ÐHÈÈÐQ€ r   )r   r#   r|   rÑ   r,   r   Úshared_memory_descriptor_typers   Úcreate_local_allocrº   r€   Úshared_memory_descriptor)ri   r[   rD   r\   r+   r‚   r€   s    ```   r   Úallocate_sharedzGluonSemantic.allocate_sharedú   sÂ   ú€ ÜŒz˜*¤d§j¡jÓ1Ó3vÔwÜŒ|˜EÓ"Ó$gÔhÜŒz˜&¤$×"3Ñ"3Ó4ÛQô	Sä×/Ñ/°
¸EÀ6È5ÓQˆØÐØ—\‘\×4Ñ4°R·X±X¸d¿l¹lÓ5KÈUÏ\É\ÓZ‰Fà—\‘\×4Ñ4°R·X±X¸d¿l¹lÓ5KÓLˆFÜ×,Ñ,¨V°ZÀÈÐPUÓVÐVr   c                 óR  ‡— t        t        ‰t        j                  «      ˆfd„«       t        j                  |j
                  |j                  ‰«      }| j                  j                  |j                  | j                  «      |j                  «      }t        j                  ||«      S )Nc                  ó   •— d‰ › �S rÚ   r   rœ   s   €r   r/   z+GluonSemantic.shared_load.<locals>.<lambda>  rÛ   r   )r   r#   r|   r   r}   rÑ   rD   rs   Úcreate_local_loadrº   r€   r   )ri   Úmem_descr\   r¼   r€   s     `  r   Úshared_loadzGluonSemantic.shared_load  sr   ø€ ÜŒz˜&¤$×"8Ñ"8Ó9ÛVô	Xä×&Ñ& x§~¡~°x·~±~ÀvÓNˆØ—‘×/Ñ/°·±¸T¿\¹\Ó0JÈHÏOÉOÓ\ˆÜ�{‰{˜6 6Ó*Ð*r   c                 óR  ‡‡— t        t        ‰t        j                  «      ˆfd„«       t        ‰j                  ‰j                  k(  ˆˆfd„«       t        ‰j
                  ‰j
                  k(  ˆˆfd„«       | j                  j                  ‰j                  ‰j                  «       y )Nc                  ó    •— dt        ‰ «      › �S )Nz+expected 'value' to be a tensor, but got a r:   r*   s   €r   r/   z,GluonSemantic.shared_store.<locals>.<lambda>  s   ø€ Ð9dÔeiÐjoÓepÐdqÐ7r€ r   c                  ó>   •— d‰j                   › d‰ j                   › d�S )Nzsource shape z and destination shape ú must matchrC   ©ró   r+   s   €€r   r/   z,GluonSemantic.shared_store.<locals>.<lambda>  ó!   ø€ ˜ u§{¡{ mÐ3JÈ8Ï>É>ÐJZÐZeÐf€ r   c                  ó>   •— d‰j                   › d‰ j                   › d�S )Nzsource dtype z and destination dtype rø   ©rÑ   rù   s   €€r   r/   z,GluonSemantic.shared_store.<locals>.<lambda>  rú   r   )	r   r#   r|   r   rD   rÑ   rs   Úcreate_local_storer€   )ri   ró   r+   s    ``r   Úshared_storezGluonSemantic.shared_store  sk   ù€ ÜŒz˜%¤§¡Ó-Ó/rÔsÜˆu�{‰{˜hŸn™nÑ,Üfô	häˆu�{‰{˜hŸn™nÑ,Üfô	hà�‰×'Ñ'¨¯©¸¿¹ÕFr   c                 óÎ  ‡‡‡— t        t        ‰t        j                  «      ˆfd„«       t        t        ‰t        «      ˆfd„«       t        t        ‰j                  «      ‰j                  k(  ˆˆfd„«       t        d‰cxk  xr ‰j                  k  nc ˆˆfd„«       t        ‰j                  j                  «       ˆfd„«       t        j                  ‰j                  ‰j                  ‰j                  j                  «      }| j                  j                  |j                  | j                  «      ‰j                   ‰j                   ‰«      }t        j                  ||«      S )Nc                  ó    •— dt        ‰ «      › �S ©Nz-expected 'indices' to be a tensor, but got a r:   ©Úindicess   €r   r/   z-GluonSemantic.shared_gather.<locals>.<lambda>  ó   ø€ ÐFÄtÈGÃ}ÀoÐV€ r   c                  ó    •— dt        ‰ «      › �S ©Nz(expected 'axis' to be an int, but got a r:   ©r”   s   €r   r/   z-GluonSemantic.shared_gather.<locals>.<lambda>  ó   ø€ Ð0XÔY]Ð^bÓYcÐXdÐ.e€ r   c                  óN   •— dt        ‰ j                  «      › d‰j                  › �S ©Nz*indices rank must match memdesc rank: got r�   ©rU   rD   r^   ©r  ró   s   €€r   r/   z-GluonSemantic.shared_gather.<locals>.<lambda>  ó(   ø€ Ð@ÄÀWÇ]Á]ÓASÐ@TÐTYÐZb×ZgÑZgÐYhÐi€ r   r   c                  ó(   •— d‰ › d‰j                   › �S ©Nzaxis z# is out of bounds for memdesc rank ©r^   ©r”   ró   s   €€r   r/   z-GluonSemantic.shared_gather.<locals>.<lambda>  ó   ø€ °E¸$¸Ð?bÐck×cpÑcpÐbqÐ2r€ r   c                  ó"   •— d‰ j                   › �S ©Nz%indices must have integer dtype, got rü   r  s   €r   r/   z-GluonSemantic.shared_gather.<locals>.<lambda>  ó   ø€ Ð1VÐW^×WdÑWdÐVeÐ/f€ r   )r   r#   r|   r   r$   rU   rD   r^   rÑ   Úis_intr}   r;   r\   rs   Úcreate_local_gatherrº   r€   )ri   ró   r  r”   r¼   r€   s    ```  r   Úshared_gatherzGluonSemantic.shared_gather  sì   ú€ ÜŒz˜'¤4§;¡;Ó/ÛVô	XäŒz˜$¤Ó$Ó&eÔfÜÜ�—‘Ó (§-¡-Ñ/Üiô	kô 	ˆq�DÖ(˜8Ÿ=™=Ô(Ô*rÔsÜˆw�}‰}×#Ñ#Ó%Ó'fÔgä×&Ñ& x§~¡~°w·}±}ÀgÇlÁl×FYÑFYÓZˆØ—‘×1Ñ1°&·,±,¸t¿|¹|Ó2LÈhÏoÉoÐ_f×_mÑ_mÐosÓtˆÜ�{‰{˜6 6Ó*Ð*r   c                 ód  ‡‡‡‡— t        t        ‰t        j                  «      ˆfd„«       t        t        ‰t        «      ˆfd„«       t        t        ‰t        j                  «      ˆfd„«       t        t        ‰j                  «      ‰j                  k(  ˆˆfd„«       t        d‰cxk  xr ‰j                  k  nc ˆˆfd„«       t        ‰j                  j                  «       ˆfd„«       t        ‰j                  ‰j                  k(  ˆˆfd„«       t        ‰j                  j                  ‰j                  j                  k(  d	„ «       t        ‰j                  ‰j                  k(  ˆˆfd
„«       | j                  j                  ‰j                  ‰j                  ‰j                  ‰«       y )Nc                  ó    •— dt        ‰ «      › �S r  r:   r  s   €r   r/   z.GluonSemantic.shared_scatter.<locals>.<lambda>%  r  r   c                  ó    •— dt        ‰ «      › �S r  r:   r  s   €r   r/   z.GluonSemantic.shared_scatter.<locals>.<lambda>&  r  r   c                  ó    •— dt        ‰ «      › �S )Nz,expected 'values' to be a tensor, but got a r:   )Úvaluess   €r   r/   z.GluonSemantic.shared_scatter.<locals>.<lambda>'  s   ø€ Ð:fÔgkÐlrÓgsÐftÐ8u€ r   c                  óN   •— dt        ‰ j                  «      › d‰j                  › �S r
  r  r  s   €€r   r/   z.GluonSemantic.shared_scatter.<locals>.<lambda>*  r  r   r   c                  ó(   •— d‰ › d‰j                   › �S r  r  r  s   €€r   r/   z.GluonSemantic.shared_scatter.<locals>.<lambda>+  r  r   c                  ó"   •— d‰ j                   › �S r  rü   r  s   €r   r/   z.GluonSemantic.shared_scatter.<locals>.<lambda>,  r  r   c                  ó<   •— d‰j                   › d‰ j                   › �S )Nz0values must have the same shape as indices: got r�   rC   )r  r  s   €€r   r/   z.GluonSemantic.shared_scatter.<locals>.<lambda>.  s"   ø€ ÐIÈ&Ï,É,ÈÐW\Ð]d×]jÑ]jÐ\kÐl€ r   c                   ó   — y)Nz+values must have the same layout as indicesr   r   r   r   r/   z.GluonSemantic.shared_scatter.<locals>.<lambda>/  r0   r   c                  ó<   •— d‰j                   › d‰ j                   › �S )Nz=values element type must match destination element type: got r�   rü   )ró   r  s   €€r   r/   z.GluonSemantic.shared_scatter.<locals>.<lambda>2  s&   ø€ ÐSÐTZ×T`ÑT`ÐSaÐafÐgo×guÑguÐfvÐw€ r   )r   r#   r|   r   r$   rU   rD   r^   rÑ   r  r;   r\   rs   Úcreate_local_scatterr€   )ri   ró   r  r  r”   s    ````r   Úshared_scatterzGluonSemantic.shared_scatter#  s  û€ ÜŒz˜'¤4§;¡;Ó/ÛVô	XäŒz˜$¤Ó$Ó&eÔfÜŒz˜&¤$§+¡+Ó.Ó0uÔvÜÜ�—‘Ó (§-¡-Ñ/Üiô	kô 	ˆq�DÖ(˜8Ÿ=™=Ô(Ô*rÔsÜˆw�}‰}×#Ñ#Ó%Ó'fÔgÜˆv�|‰|˜wŸ}™}Ñ,Ülô	näˆv�{‰{×!Ñ! W§\¡\×%8Ñ%8Ñ8Ñ:oÔpÜØ�L‰L˜HŸN™NÑ*Üwô	yð 	�‰×)Ñ)¨(¯/©/¸6¿=¹=È'Ï.É.ÐZ^Õ_r   c                 óŠ  — t        |t        j                  «      st        dt	        |«      › �«      ‚t        |t        j
                  «      st        dt	        |«      › �«      ‚|j                  |j                  k7  r&t        d|j                  › d|j                  › d�«      ‚|j                  |j                  k7  r&t        d|j                  › d|j                  › d�«      ‚|j                  |j                  t        |j                  «       d  k7  r%t        d	|j                  › d
|j                  › �«      ‚|j                  j                  | j                  «      }|j                  j                  | j                  «      }| j                  j                  ||t        |j                  «      |j                  j                   «      S )NzIbank_conflicts expects the register layout to be a distributed_type, got zTbank_conflicts expects the shared layout to be a shared_memory_descriptor_type, got zregister shape z and shared shape rø   z$mismatched dtypes between register (z) and shared (z	) layoutsz,bank_conflicts NYI for subslices. Got shape z and alloc_shape )r#   r|   r}   rÞ   r;   rì   rD   rŽ   r[   Úalloc_shaperU   r\   Ú_to_irrs   Úget_shared_bank_conflictsrT   rX   )ri   Údistr_tyÚ	shared_tyÚreg_attrÚshared_attrs        r   Úbank_conflictszGluonSemantic.bank_conflicts6  s£  € Ü˜(¤D×$9Ñ$9Ô:ÜØ[Ô\`ÐaiÓ\jÐ[kÐlónð nô ˜)¤T×%GÑ%GÔHÜØfÔgkÐluÓgvÐfwÐxóð ð �>‰>˜YŸ_™_Ò,Ü˜¨x¯~©~Ð.>Ð>PÐQZ×Q`ÑQ`ÐPaÐalÐmÓnÐnØ×Ñ 8×#6Ñ#6Ò6ÜØ6°x×7JÑ7JÐ6KÈ>ÐZc×ZnÑZnÐYoÐoxÐyóð ð �?‰?˜i×3Ñ3´S¸¿¹Ó5IÐ4IÐ4JÐKÒKÜØ>¸y¿¹Ð>OÐO`Ðaj×avÑavÐ`wÐxóð ð —?‘?×)Ñ)¨$¯,©,Ó7ˆØ×&Ñ&×-Ñ-¨d¯l©lÓ;ˆØ�|‰|×5Ñ5°hÀÌTÐRZ×R`ÑR`ÓMaØ6>×6IÑ6I×6\Ñ6\ó^ð 	^r   c           	      óN  ‡— ddl m}m} t        t	        ‰t
        t        ||f«      ˆfd„«       t	        ‰t        t        t        f«      rt        j                  ‰«      S t        j                  | j                  j                  ‰j                  | j                  «      |«      «      S )Nr   )ÚTensorMemoryLayoutÚTensorMemoryScalesLayoutc                  ó    •— dt        ‰ «      › �S )NzcExpected a DistributedLayout, SharedLayout, or TensorMemoryLayout or TensorMemoryScalesLayout, got r:   rœ   s   €r   r/   z0GluonSemantic.to_linear_layout.<locals>.<lambda>V  s   ø€ ØqÔrvÐw}Ór~Ðqð  Að r   )Ú3triton.experimental.gluon.language.nvidia.blackwellr0  r1  r   r#   r   r   r   r   r   r|   Ú	constexprrs   rß   r(  )ri   r\   rD   r0  r1  s    `   r   rß   zGluonSemantic.to_linear_layoutP  s�   ø€ ÷	
ô 	Ü�vÔ 1´<ÐASÐUmÐnÓoó rAô	
ô
 �fœzÔ+BÔDVÐWÔXÜ—>‘> &Ó)Ð)ä�~‰~˜dŸl™l×;Ñ;¸F¿M¹MÈ$Ï,É,Ó<WÐY^Ó_Ó`Ð`r   c                 óN   — | j                   j                  |j                  «       y r   )rs   Úcreate_local_deallocr€   )ri   ró   s     r   Úshared_dealloczGluonSemantic.shared_dealloc_  s   € Ø�‰×)Ñ)¨(¯/©/Õ:r   c                 óž  ‡‡— ‰j                   }t        t        ‰t        «      ˆfd„«       t        t        |j                  t
        «      ˆfd„«       | j                  j                  ‰j                  | j                  «      ‰j                  «      }t        j                  |j                  |j                  ‰«      }| j                  ||«      S )Nc                  ó   •— d‰ › �S )Nz9set_auto_layout must set to a distributed layout but got r   rœ   s   €r   r/   z/GluonSemantic.set_auto_layout.<locals>.<lambda>e  s   ø€ ÐRÐSYÐRZÐ[€ r   c                  ó6   •— d‰ j                   j                  › �S )Nz4set_auto_layout input must have auto layout but got )r;   r\   r*   s   €r   r/   z/GluonSemantic.set_auto_layout.<locals>.<lambda>g  s   ø€ ÐMÈeÏjÉj×N_ÑN_ÐM`Ða€ r   )r;   r   r#   r   r\   r   rs   Úcreate_set_auto_layoutr(  r€   r|   r}   r[   rD   r   )ri   r+   r\   Úsrc_tyr€   Úres_tys    ``   r   rÄ   zGluonSemantic.set_auto_layoutb  s”   ù€ Ø—‘ˆÜŒz˜&Ô"3Ó4Û[ô	]äŒz˜&Ÿ-™-¬Ó4Ûaô	cà—‘×4Ñ4°V·]±]À4Ç<Á<Ó5PÐRW×R^ÑR^Ó_ˆÜ×&Ñ& v×'8Ñ'8¸&¿,¹,ÈÓOˆØ�{‰{˜6 6Ó*Ð*r   c                 óH  ‡‡‡— t        t        ‰t        «      ˆfd„«       t        t        ‰t        «      ˆfd„«       t        t        ‰t        «      ˆfd„«       dg|j                  z  }‰|‰<   t	        |j
                  «      }‰|‰<   |j                  }t        j                  |j                  |||j                  j                  «      }| j                  }	|	j                  |j                  |	«      |j                  |«      }
t        j                   |
fi |j"                  ¤ŽS )Nc                  ó   •— d‰ › �S )Nz&expected 'start' to be an int but got r   )rÊ   s   €r   r/   z-GluonSemantic.memdesc_slice.<locals>.<lambda>m  s   ø€ Ð1WÐX]ÐW^Ð/_€ r   c                  ó   •— d‰ › �S )Nz'expected 'length' to be an int but got r   )Úlengths   €r   r/   z-GluonSemantic.memdesc_slice.<locals>.<lambda>n  s   ø€ Ð2YÐZ`ÐYaÐ0b€ r   c                  ó   •— d‰ › �S )Nz$expected 'dim' to be an int but got r   rž   s   €r   r/   z-GluonSemantic.memdesc_slice.<locals>.<lambda>o  s   ø€ Ð/SÐTWÐSXÐ-Y€ r   r   )r   r#   r$   r^   rT   rD   r\   r|   rì   rÑ   r;   r'  rs   Úcreate_memdesc_subslicerº   r€   rî   Ú__dict__)ri   ró   rÊ   rA  rA   ÚoffsetsrD   r\   r‚   rs   r€   s     ```      r   Úmemdesc_slicezGluonSemantic.memdesc_slicel  sÛ   ú€ ÜŒz˜%¤Ó%Ó'_Ô`ÜŒz˜&¤#Ó&Ó(bÔcÜŒz˜#œsÓ#Ó%YÔZØ�#˜Ÿ™Ñ%ˆØˆ�‰Ü�X—^‘^Ó$ˆØˆˆc‰
Ø—‘ˆÜ×/Ñ/°·±ÀÀvÈxÏ}É}×OhÑOhÓiˆØ—,‘,ˆØ×0Ñ0°·±¸'Ó1BÀHÇOÁOÐU\Ó]ˆÜ×,Ñ,¨VÑC°r·{±{ÑCÐCr   c                 óÚ  ‡— | j                  ‰«      Št        ‰j                  t        j                  k(  ˆfd„«       |j
                  dd  }| j                  ‰«      j                  Š|j                  }t        j                  |j                  |||«      }| j                  }|j                  |j                  |«      |j                  ‰«      }t        j                  |fi |j                  ¤ŽS )Nc                  ó"   •— d‰ j                   › �S )Nz%expected 'index' to be int32 but got r:   ©Úindexs   €r   r/   z-GluonSemantic.memdesc_index.<locals>.<lambda>|  s   ø€ Ð3XÐY^×YcÑYcÐXdÐ1e€ r   r	   )Ú	to_tensorr   r;   r|   rÈ   rD   r€   r\   rì   rÑ   rs   Úcreate_memdesc_indexrº   rî   rD  )ri   ró   rJ  rD   r\   r‚   rs   r€   s     `     r   Úmemdesc_indexzGluonSemantic.memdesc_indexz  sµ   ø€ Ø—‘˜uÓ%ˆÜˆu�z‰zœTŸZ™ZÑ'Ó)eÔfØ—‘˜q˜rÐ"ˆØ—‘˜uÓ%×,Ñ,ˆØ—‘ˆÜ×/Ñ/°·±ÀÀvÈuÓUˆØ—,‘,ˆØ×-Ñ-¨b¯h©h°wÓ.?ÀÇÁÐRWÓXˆÜ×,Ñ,¨VÑC°r·{±{ÑCÐCr   c                 ój  ‡‡— t        t        ‰«      ˆfd„«       t        t        ‰«      t        ‰j                  «      k(  ˆˆfd„«       ‰D �cg c]  }‰j                  |   ‘Œ }}‰j                  j
                  }|d t        |«      ‰j                  z
   }|‰D �cg c]   }|t        |«      ‰j                  z
  d  |   ‘Œ" c}z  }| j                  j                  ‰j                  ‰«      }| j                  j                  |«      }t        j                  |‰j                  |||¬«      S c c}w c c}w )Nc                  ó   •— d‰ › �S )Nz1all elements of 'order' must be integers but got r   )Úorders   €r   r/   z-GluonSemantic.memdesc_trans.<locals>.<lambda>†  rè   r   c                  ó<   •— d‰ j                   › dt        ‰«      › d�S )Nzsource rank (z) and order length (z) must match)r^   rU   )ró   rP  s   €€r   r/   z-GluonSemantic.memdesc_trans.<locals>.<lambda>‰  s    ø€ �m H§M¡M ?Ð2FÄsÈ5ÃzÀlÐR^Ð_€ r   ©r[   rD   r'  r\   )r   r,   rU   rD   r;   r'  r^   rs   Úcreate_memdesc_transr€   Úget_gluon_layout_from_memdescr|   rî   rÑ   )	ri   ró   rP  r&   rD   r'  Únew_alloc_shaper€   r\   s	    ``      r   Úmemdesc_transzGluonSemantic.memdesc_trans…  s  ù€ ÜŒ|˜EÓ"Ó$gÔhÜÜ�‹Jœ#˜hŸn™nÓ-Ñ-Ü_ô	að -2Ö2 q�—‘ Ó"Ð2ˆÐ2Ø—m‘m×/Ñ/ˆØ%Ð&G¤s¨;Ó'7¸(¿-¹-Ñ'GÐHˆØÐW\Ö]ÐRS˜K¬¨KÓ(8¸8¿=¹=Ñ(HÐ(IÐJÈ1ÓMÒ]Ñ]ˆà—‘×2Ñ2°8·?±?ÀEÓJˆØ—‘×;Ñ;¸FÓCˆÜ×,Ñ,¨VÀÇÁÐV[Ø9HÐQWôYð 	Yùò 3ùò ^s   ÁD+Â%D0c                 ó  ‡‡— t        t        ‰«      ˆfd„«       t        t        j                  ‰«      t        j                  ‰j                  «      k(  ˆˆfd„«       | j
                  j                  ‰j                  ‰«      }| j
                  j                  |«      }‰j                  j                  }t        |«      ‰j                  z
  }|d | t        ‰«      z   }t        j                  |‰j                   ‰||¬«      S )Nc                  ó   •— d‰ › �S rç   r   rC   s   €r   r/   z/GluonSemantic.memdesc_reshape.<locals>.<lambda>–  rè   r   c                  ó(   •— d‰ j                   › d‰› �S )Nz)memdesc_reshape total elements mismatch: z -> rC   )ró   rD   s   €€r   r/   z/GluonSemantic.memdesc_reshape.<locals>.<lambda>™  s   ø€ Ð@Ø Ÿ™Ð' t¨E¨7ð4€ r   rR  )r   r,   ÚmathÚprodrD   rs   Úcreate_memdesc_reshaper€   rT  r;   r'  rU   r^   rT   r|   rî   rÑ   )ri   ró   rD   r€   r\   r'  Ú
prefix_lenrU  s    ``     r   Úmemdesc_reshapezGluonSemantic.memdesc_reshape•  sË   ù€ ÜŒ|˜EÓ"Ó$gÔhÜÜ�I‰I�eÓ¤§	¡	¨(¯.©.Ó 9Ñ9ô5ô	
ð —‘×4Ñ4°X·_±_ÀeÓLˆØ—‘×;Ñ;¸FÓCˆØ—m‘m×/Ñ/ˆÜ˜Ó%¨¯©Ñ5ˆ
Ø% k zÐ2´T¸%³[Ñ@ˆä×,Ñ,ØØ—~‘~ØØ'Øô
ð 	
r   c                 óÂ  ‡‡‡— t        t        ‰t        j                  «      ˆfd„«       t        t	        ‰«      ˆfd„«       t        t        ‰t        j
                  «      ˆfd„«       t        j                  ‰‰‰‰«      }| j                  j                  |j                  | j                  «      |j                  «      }t        j                  |fi |j                  ¤ŽS )Nc                  ó   •— d‰ › �S )Nz'expected 'dtype' to be a dtype but got r   rü   s   €r   r/   z3GluonSemantic.memdesc_reinterpret.<locals>.<lambda>¬  s   ø€ Ð8_Ð`eÐ_fÐ6g€ r   c                  ó   •— d‰ › �S rç   r   rC   s   €r   r/   z3GluonSemantic.memdesc_reinterpret.<locals>.<lambda>­  rè   r   c                  ó   •— d‰ › �S rê   r   rœ   s   €r   r/   z3GluonSemantic.memdesc_reinterpret.<locals>.<lambda>¯  rë   r   )r   r#   r|   rÑ   r,   r   rì   rs   Úcreate_memdesc_reinterpretrº   r€   rî   rD  )ri   ró   rÑ   rD   r\   r‚   r€   s     ```  r   Úmemdesc_reinterpretz!GluonSemantic.memdesc_reinterpret«  sž   ú€ ÜŒz˜%¤§¡Ó,Ó.gÔhÜŒ|˜EÓ"Ó$gÔhÜŒz˜&¤$×"3Ñ"3Ó4ÛQô	Sä×/Ñ/°°u¸fÀeÓLˆØ—‘×8Ñ8¸¿¹À$Ç,Á,Ó9OÐQY×Q`ÑQ`ÓaˆÜ×,Ñ,¨VÑC°r·{±{ÑCÐCr   c                 ó^   — |rt        j                  |||«      }n|}| j                  ||«      S r   )r|   r}   r   )ri   r¢   r�   r�   r\   r=  s         r   Úwrap_tensorzGluonSemantic.wrap_tensor´  s/   € ÙÜ×*Ñ*¨9°iÀÓH‰FàˆFØ�{‰{˜1˜fÓ%Ð%r   c                 ó  ‡‡‡— | D ]3  Št        t        ‰j                  t        j                  «      ˆfd„«       Œ5 | D �cg c]  }|j                  j
                  ‘Œ c}Š‰d   Št        t        ˆfd„‰dd  D «       «      ˆfd„«       y c c}w )Nc                  ó"   •— d‰ j                   ›�S ©Nz#expected distributed_type but got: r:   )r¢   s   €r   r/   z2GluonSemantic._check_same_layout.<locals>.<lambda>¾  s   ø€ ÐHkÐlm×lrÑlrÐkuÐFv€ r   r   c              3   ó(   •K  — | ]	  }|‰k(  –— Œ y ­wr   r   )r%   ÚlÚl0s     €r   r'   z3GluonSemantic._check_same_layout.<locals>.<genexpr>Á  s   øè ø€ Ò0˜q�1˜•7Ñ0ùs   ƒr	   c                  ó   •— d‰ › �S )Nz3Expected inputs to have matching layouts, but got: r   )Úlayoutss   €r   r/   z2GluonSemantic._check_same_layout.<locals>.<lambda>Â  s   ø€ ÐLÈWÈIÐV€ r   )r   r#   r;   r|   r}   r\   r)   )Úxsr¢   rl  rn  s    `@@r   Ú_check_same_layoutz GluonSemantic._check_same_layout»  su   ú€ àò 	xˆAÜ”:˜aŸf™f¤d×&;Ñ&;Ó<Ó>vÕwð	xà*,Ö- Q�1—6‘6—=“=Ò-ˆØ�Q‰ZˆÜŒsÓ0 G¨A¨B KÔ0Ó0ÛVõ	Xùò .s   Á B
ÚinputsÚreverse.c                 óæ  ‡ ‡‡‡— ‰d   j                   j                  Št        ‰«      }| |cxk  r|k  sn J d|› d|› d�«       ‚|dk  r||z  }‰D ]"  }|j                   j                  ‰k(  rŒJ d«       ‚ ‰ j                  j	                  ‰D �cg c]  }|j
                  ‘Œ c}||«      Š |‰«       ‰j                  «       sJ ‚t        ˆˆˆ ˆfd„t        t        ‰«      «      D «       «      S c c}w )Nr   z
scan axis z must be < inputs rank (ú)z(all scan inputs must have the same shapec              3   ó�   •K  — | ]=  }‰j                  ‰j                  |«      ‰|   j                  j                  ‰«      –— Œ? y ­wr   ©rƒ   Ú
get_resultr;   r…   )r%   r&   rq  Úscan_opri   rD   s     €€€€r   r'   z1GluonSemantic.associative_scan.<locals>.<genexpr>Õ  sB   øè ø€ ò )àð ×*Ñ*¨7×+=Ñ+=¸aÓ+@À&ÈÁ)Ç.Á.×BWÑBWÐY^×_ñ)ùó   ƒAA)	r;   rD   rU   rs   Úcreate_scanr€   ÚverifyÚtupleÚrange)	ri   rq  r”   Úregion_builder_fnrr  r^   Útrx  rD   s	   ``     @@r   Úassociative_scanzGluonSemantic.associative_scanÄ  sò   û€ à�q‘	—‘×$Ñ$ˆÜ�5‹zˆàˆu˜Ô#˜tÔ#ÐW z°$°Ð7OÐPTÈvÐUVÐ%WÓWÐ#à�!Š8Ø�D‰LˆDàò 	UˆAØ—6‘6—<‘< 5Ó(ÐTÐ*TÓTÐ(ð	Uð —,‘,×*Ñ*¸fÖ+E¸¨A¯H«HÒ+EÀtÈWÓUˆÙ˜'Ô"Ø�~‰~ÔÐÐäö )äœ3˜v›;Ó'ô)ó )ð 	)ùò	 ,Fs   ÂC.c                 óh  ‡ ‡‡‡‡‡	‡
— ‰€t        ˆ fd„‰D «       «      ŠdŠ‰d   j                  j                  Š
t        ‰
«      Št	        d‰cxk  xr ‰k  nc ˆˆfd„«       ‰ j                  ‰«       t        ‰
«      D ��cg c]  \  }}|‰k7  sŒ|‘Œ c}}Š	t        ˆ
fd„‰D «       «      sJ d«       ‚‰ j                  j                  ‰D �cg c]  }|j                  ‘Œ c}‰«      Š |‰«       ‰j                  «       sJ ‚t        ˆˆˆ	ˆ fd„t        t        ‰«      «      D «       «      S c c}}w c c}w )Nc              3   óp   •K  — | ]-  }‰j                  ||j                  j                  gd ¬«      –— Œ/ y­w)F)rÌ   N)rÏ   Únumelr+   )r%   r  ri   s     €r   r'   z*GluonSemantic.reduction.<locals>.<genexpr>Û  s+   øè ø€ Ò_ÐST˜4Ÿ<™<¨¨A¯G©G¯M©M¨?È˜<×NÑ_ùs   ƒ36r   c                  ó   •— d‰› d‰ › �S )Nz/expected reduction axis to be in the range [0, z
) but got r   ©r”   r^   s   €€r   r/   z)GluonSemantic.reduction.<locals>.<lambda>à  s   ø€ Ð+ZÐ[_ÐZ`Ð`jÐkoÐjpÐ)q€ r   c              3   óP   •K  — | ]  }|j                   j                  ‰k(  –— Œ y ­wr   )r;   rD   )r%   r  rD   s     €r   r'   z*GluonSemantic.reduction.<locals>.<genexpr>ã  s   øè ø€ Ò9¨Q�1—6‘6—<‘< 5Õ(Ñ9ùs   ƒ#&z-all reduction inputs must have the same shapec              3   ó�   •K  — | ]=  }‰j                  ‰j                  |«      ‰|   j                  j                  ‰«      –— Œ? y ­wr   rv  )r%   r&   rq  Ú	reduce_opr�   ri   s     €€€€r   r'   z*GluonSemantic.reduction.<locals>.<genexpr>é  sB   øè ø€ ò )àð ×*Ñ*¨9×+?Ñ+?ÀÓ+BÀFÈ1ÁIÇNÁN×DYÑDYÐ[d×eñ)ùry  )r|  r;   rD   rU   r   rp  rZ   r)   rs   Úcreate_reducer€   r{  r}  )ri   rq  r”   r~  r&   Úsr  r^   rˆ  r�   rD   s   ```    @@@@r   Ú	reductionzGluonSemantic.reductionÙ  s  þ€ Øˆ<ÜÓ_ÐX^Ô_Ó_ˆFØˆDà�q‘	—‘×$Ñ$ˆÜ�5‹zˆÜˆq�DÖ˜4ÔÔ!qÔrØ×Ñ Ô'Ü#,¨UÓ#3×A™4˜1˜a°q¸D³y’QÓAˆ	ÜÓ9°&Ô9Ô9ÐjÐ;jÓjÐ9à—L‘L×.Ñ.À&Ö/I¸Q°·³Ò/IÈ4ÓPˆ	Ù˜)Ô$Ø×ÑÔ!Ð!Ð!äö )äœ3˜v›;Ó'ô)ó )ð 	)ùó Bùò 0Js   ÂD)ÂD)ÃD/Únum_binsÚmaskc                 ó   — t        t        |j                  «      dk(  d„ «       t        |j                  j	                  «       d„ «       t        |d ud„ «       |�P| j                  ||«      \  }}t        |j                  j                  j                  «       d„ «       |j                  }|j                  | j                  «      }| j                  j                  |j                  |||«      }| j                  |t        j                  |g|«      S )Nr	   c                   ó   — y)Nz histogram only supports 1D inputr   r   r   r   r/   z)GluonSemantic.histogram.<locals>.<lambda>î  r0   r   c                   ó   — y)Nz%histogram only supports integer inputr   r   r   r   r/   z)GluonSemantic.histogram.<locals>.<lambda>ï  r0   r   c                   ó   — y)Nz'histogram requires a destination layoutr   r   r   r   r/   z)GluonSemantic.histogram.<locals>.<lambda>ð  r0   r   c                   ó   — y)Nz"Mask must have boolean scalar typer   r   r   r   r/   z)GluonSemantic.histogram.<locals>.<lambda>ó  r0   r   )r   rU   rD   rÑ   r  r©   r;   r…   Úis_boolr€   r(  rs   Úcreate_histogramrf  r|   rÈ   )ri   r“   rŒ  r�  r\   Úlayout_attrr€   s          r   Ú	histogramzGluonSemantic.histogramí  sÓ   € ÜŒs�5—;‘;Ó 1Ñ$Ñ&PÔQÜˆu�{‰{×!Ñ!Ó#Ñ%TÔUÜˆv˜TÐ!Ñ#TÔUØÐØ×3Ñ3°D¸%Ó@‰KˆD�%Ü�4—9‘9×#Ñ#×+Ñ+Ó-Ñ/[Ô\Ø—;‘;ˆDØ—m‘m D§L¡LÓ1ˆØ—‘×.Ñ.¨u¯|©|¸XÀtÈ[ÓYˆØ×Ñ ¬¯
©
°X°JÀÓGÐGr   c           	      óÆ  — t        |d ud„ «       t        |d„ «       t        t        |j                  «      dk(  d„ «       t        j                  |j
                  j                  |j                  d   |j                  d   z   g|«      }| j                  | j                  j                  |j                  |j                  |j                  | j                  «      «      |«      S )Nc                   ó   — y)Nz!cat requires a destination layoutr   r   r   r   r/   z#GluonSemantic.cat.<locals>.<lambda>ú  r0   r   c                   ó   — y)Nz;current implementation of `cat` always may reorder elementsr   r   r   r   r/   z#GluonSemantic.cat.<locals>.<lambda>û  r0   r   r	   c                   ó   — y)Nzcat requires a rank-1 inputr   r   r   r   r/   z#GluonSemantic.cat.<locals>.<lambda>ü  r0   r   r   )r   rU   rD   r|   r}   r;   r…   r   rs   Ú
create_catr€   rº   )ri   r¯   r°   rÌ   r\   Úret_types         r   ÚcatzGluonSemantic.catù  s¨   € Üˆv˜TÐ!Ñ#NÔOÜˆ{ÑaÔbÜŒs�3—9‘9‹~ Ñ"Ñ$IÔJÜ×(Ñ(¨¯©¯©¸3¿9¹9ÀQ¹<È#Ï)É)ÐTUÉ,Ñ;VÐ:WÐY_Ó`ˆØ�{‰{˜4Ÿ<™<×2Ñ2°3·:±:¸s¿z¹zÈ8Ï>É>ÐZ^×ZfÑZfÓKgÓhÐjrÓsÐsr   ÚsrcrJ  c                 óÈ  ‡‡‡‡— t        t        ‰j                  t        j                  «      ˆfd„«       t        t        ‰j                  t        j                  «      ˆfd„«       t        ‰j                  j
                  j                  «       ˆfd„«       t        ‰j                  j                  «      Št        t        ‰j                  j                  «      ‰k(  d„ «       t        ‰ ‰cxk  xr ‰k  nc ˆˆfd„«       ‰dk  r‰‰z  Št        ‰«      D ]H  }|‰k(  rŒ	t        ‰j                  j                  |   ‰j                  j                  |   k(  ˆfd„«       ŒJ | j                  j                  ‰j                  ‰j                  ‰«      }| j                  |‰j                  j
                  ‰j                  j                  ‰j                  j                  «      S )Nc                  ó"   •— d‰ j                   ›�S ri  r:   )rž  s   €r   r/   z&GluonSemantic.gather.<locals>.<lambda>  s   ø€ ÐFiÐjm×jrÑjrÐiuÐDv€ r   c                  ó"   •— d‰ j                   ›�S ri  r:   rI  s   €r   r/   z&GluonSemantic.gather.<locals>.<lambda>  s   ø€ Ð<¸U¿Z¹Z¸NÐK€ r   c                  ó6   •— d‰ j                   j                  ›�S )Nz&expected integer scalar type but got: )r;   r…   rI  s   €r   r/   z&GluonSemantic.gather.<locals>.<lambda>  s   ø€ Ð5[Ð\a×\fÑ\f×\mÑ\mÐ[pÐ3q€ r   c                   ó   — y)Nz0source and index tensors must have the same rankr   r   r   r   r/   z&GluonSemantic.gather.<locals>.<lambda>  r0   r   c                  ó   •— d‰ › d‰› d�S )Nzgather axis z must be < source rank (rt  r   r…  s   €€r   r/   z&GluonSemantic.gather.<locals>.<lambda>  s   ø€ ¨|¸D¸6ÐAYÐZ^ÐY_Ð_`Ð-a€ r   r   c                  ó   •— d‰ › d�S )Nz
index dim z( must match the corresponding source dimr   r  s   €r   r/   z&GluonSemantic.gather.<locals>.<lambda>  s   ø€ ˜* T FÐ*RÐS€ r   )r   r#   r;   r|   r}   r…   r  rU   rD   r}  rs   Úcreate_gatherr€   rf  r\   )ri   rž  rJ  r”   ÚdÚgatherr^   s    ```  @r   r¨  zGluonSemantic.gather   sU  û€ ÜŒz˜#Ÿ(™(¤D×$9Ñ$9Ó:Ó<vÔwÜŒz˜%Ÿ*™*¤d×&;Ñ&;Ó<ÛKô	Mäˆu�z‰z× Ñ ×'Ñ'Ó)Ó+qÔrä�3—8‘8—>‘>Ó"ˆÜŒs�5—:‘:×#Ñ#Ó$¨Ñ,Ñ.hÔiÜ�ˆu˜Ö#˜tÔ#Ô%aÔbØ�!Š8Ø�D‰LˆDä�t“ò 	ˆAØ�DŠyØÜØ—
‘
× Ñ  Ñ# s§x¡x§~¡~°aÑ'8Ñ8ÛSõð	ð —‘×+Ñ+¨C¯J©J¸¿¹ÀdÓKˆØ×Ñ ¨¯©¯©¸¿¹×9IÑ9IÈ5Ï:É:×K\ÑK\Ó]Ð]r   c                 ó   — | j                   j                  |j                  |j                  | j                   «      |«      }t	        |j
                  j                  «      }||xx   dz  cc<   | j                  |||«      S )NrE   )rs   Úcreate_fp4_to_fpr€   rº   rT   r;   rD   rƒ   )ri   rž  Ú	elem_typer”   ÚresultrD   s         r   Ú	fp4_to_fpzGluonSemantic.fp4_to_fp  sa   € Ø—‘×.Ñ.¨s¯z©z¸9¿?¹?È4Ï<É<Ó;XÐZ^Ó_ˆÜ�S—X‘X—^‘^Ó$ˆØˆd‹�qÑ‹Ø×-Ñ-¨f°iÀÓGÐGr   Úworker_num_warpsÚworker_num_regsc                 óä  ‡— |D ]2  \  }Št        t        ‰t        t        j                  f«      ˆfd„«       Œ4 t	        |«      dk\  sJ d«       ‚|d   \  }}t	        |«      dz
  }|dd  }	|t	        |«      k(  sJ d|› dt	        |«      › d�«       ‚|�%|t	        |«      k(  sJ d|› dt	        |«      › d�«       ‚| j
                  }
|
j                  «       }|
j                  «       }|
j                  |«       |j                  ||i ¬	«      }t        |g«      }|
j                  |«       |D �cg c]  }|j                  «       ‘Œ }}|	D ��cg c]  \  }}t        |«      ‘Œ }}}t        |g «      }|
j                  |«       |
j                  ||«      }|j!                  «       j#                  |«       |�|j%                  |«       |
j'                  |j)                  «       g «       |
j+                  ||«      }|D �cg c]  }|j                  «       ‘Œ }}d}t-        |	«      D ]Ç  \  }\  }Št/        ||   ¬
«      }|
j'                  |j1                  |«      |«      }||   }t3        t	        |«      «      D �cg c]  }|j5                  ||z   «      ‘Œ }}t7        |‰D �cg c]  }|j8                  ‘Œ c}«      }|j                  ||i |¬«       |
j;                  «        |t	        |«      z  }ŒÉ |
j=                  |j?                  «       «       t3        t	        |«      «      D �cg c]  }|jA                  |«      ‘Œ }}tC        t7        ||j8                  g«      «      S c c}w c c}}w c c}w c c}w c c}w c c}w )Nc                  ó    •— dt        ‰ «      › �S )Nz9function arguments must be a tuple of arguments, but got r:   )Úargss   €r   r/   z/GluonSemantic.warp_specialize.<locals>.<lambda>   s   ø€ ÐVÔW[Ð\`ÓWaÐVbÐc€ r   r	   z8expected at least one function for the default partitionr   zwarp specialize got z partitions but z warp countsz register counts)Úkwargsr<   )r³  Úcaller_context)"r   r#   r|  r|   rU   rs   Úget_insertion_pointÚ	new_blockÚset_insertion_point_to_startÚcall_JitFunctionr   Úcreate_warp_yieldÚget_typeÚsumÚrestore_insertion_pointÚcreate_warp_specializeÚget_default_regionÚ	push_backÚset_requested_registersÚcreate_block_with_parentÚget_partition_op_holderÚ!create_warp_specialize_partitionsrZ   rg   Ú
get_regionr}  Úget_argumentr   r;   Úcreate_warp_returnÚset_insertion_point_afterÚget_operationrw  Únext)ri   Úfunctions_and_argsr®  r¯  Ú	generatorÚ_Údefault_partitionÚdefault_argsÚnum_partitionsÚworkersrs   Ú	insert_ptÚdefault_blockÚdefault_resultÚmlir_resultsÚrÚresult_typesr²  Úworker_argsÚ	mlir_argsÚws_opÚpartitions_opÚargÚ	arg_typesÚarg_itr&   Úfuncr´  ÚblockÚjÚ
block_argss                    `             r   Úwarp_specializezGluonSemantic.warp_specialize  s’  ø€ à)ò 	e‰GˆAˆtÜ”:˜d¤U¬D¯J©JÐ$7Ó8Ûcõeð	eô Ð%Ó&¨!Ò+ÐgÐ-gÓgÐ+Ø*<¸QÑ*?Ñ'Ð˜<ÜÐ/Ó0°1Ñ4ˆØ$ Q RÐ(ˆà¤Øó"
ò 
ð 	fà! .Ð!1Ð1AÄ#ÐFVÓBWÐAXÐXdÐeó	fð 
ð Ð&Ø!¤SØó&ò ð mà% nÐ%5Ð5EÄcÈ/ÓFZÐE[Ð[kÐlómð ð —,‘,ˆØ×/Ñ/Ó1ˆ	ð  ×)Ñ)Ó+ˆØ×,Ñ,¨]Ô;Ø"×3Ñ3Ð4EÀ|Ð\^Ð3Ó_ˆÜ+¨^Ð,<Ó=ˆØ×!Ñ! ,Ô/Ø.:Ö;¨˜Ÿ
™
�Ð;ˆÐ;ð BI×I±g°a¸Ô+¨DÕ1ÐIˆÑIÜ˜ RÓ(ˆ	Ø×'Ñ'¨	Ô2Ø×.Ñ.¨|Ð=MÓNˆØ× Ñ Ó"×,Ñ,¨]Ô;àÐ&Ø×)Ñ)¨/Ô:ð 	×(Ñ(¨×)FÑ)FÓ)HÈ"ÔMØ×AÑAÀ)È^Ó\ˆØ/8Ö9¨�S—\‘\•^Ð9ˆ	Ð9ØˆÜ(¨Ó1ò 	%‰OˆA‰|��dÜ/Ð:JÈ1Ñ:MÔNˆNØ×4Ñ4°]×5MÑ5MÈaÓ5PÐR[Ó\ˆEØ# A™ˆIÜBGÌÈIËÓBWÖX¸Q˜%×,Ñ,¨V°a©ZÕ8ÐXˆJÐXÜ,¨ZÈdÖ9SÀs¸#¿(»(Ò9SÓTˆJØ×&Ñ& t¨ZÀÐSaÐ&ÔbØ×&Ñ&Ô(Ø”c˜)“nÑ$‰Fð	%ð 	×)Ñ)¨%×*=Ñ*=Ó*?Ô@Ü5:¼3¸|Ó;LÓ5MÖN°˜×(Ñ(¨Õ+ÐNˆÐNÜÔ'¨°~×7JÑ7JÐ6KÓLÓMÐMùò; <ùó Jùò :ùò YùÚ9Sùò Os$   Ä'MÅMÇ4MÉ.M#ÊM(ÌM-c                 óh   — t        j                  | j                  j                  j                  «      S r   )r|   r4  rs   ÚoptionsÚnum_ctasrm   s    r   rå  zGluonSemantic.num_ctasX  s!   € Ü�~‰~˜dŸl™l×2Ñ2×;Ñ;Ó<Ð<r   c                 ó
  — |j                   �Et        |j                   t        «      sJ ‚t        j                  |j                   j
                  «      S t        j                  | j                  j                  j
                  «      S r   )r´  r#   rg   r|   r4  r=   rs   rä  )ri   rË  s     r   r=   zGluonSemantic.num_warps[  s_   € Ø×#Ñ#Ð/Ü˜i×6Ñ6Ô8JÔKÐKÐKÜ—>‘> )×":Ñ":×"DÑ"DÓEÐEÜ�~‰~˜dŸl™l×2Ñ2×<Ñ<Ó=Ð=r   )F)8ru   rv   rw   r|   r   Úlangr   Ú__annotations__rj   rƒ   r†   r   r$   r’   r   r¤   r«   r   r®   r³   r½   r©   rÉ   ÚboolrÏ   rÓ   rÖ   rã   rï   rô   rþ   r  r%  r.  rß   r7  rÄ   rF  rM  rV  r^  rd  rf  Ústaticmethodrp  r   r€  r‹  r–  r�  r¨  r­  râ  rå  r=   Ú__classcell__)r¬   s   @r   ry   ry   a   s¥  ø… Ø�[‰[€FØ€DàÓð ó ò'ò_ð¨4°©9ð ÀÀcÁó ð T ð T°ð T¸ó Tð&5�hð 5 8ð 5°õ 5ðX�xð X E¨(°HÐ*<Ñ$=õ Xð5˜Xð 5¨U°3©Zð 5¸Hõ 5ð+¨(ð +¸5À¹:ð +È(ó +ð ¨ð °xð ÀHõ ô:9ð5˜Xð 5°$°s±)ð 5È$õ 5ò
+ò1ó+ò
Wò+òGò+ò`ò&^ò4aò;ò+òDò	DòYò 
ò,Dò&ð ñXó ðXð) x°Ñ'9ð )Àð )Ø"&ð)Ø+0°¸3°Ñ+?ó)ð*) ¨Ñ 2ð )¸#ð )ÐUZÐ[cÐehÐ[hÑUió )ð(
H˜xð 
H°3ð 
H¸hð 
HÐS[ó 
Hðt�xð t hð t¸Tð tÈhó tð^˜(ð ^¨8ð ^¸3ð ^À8ó ^ð,H˜Xð H¸8ó Hð:NÀHÈSÁMð :NÐdlÐmpÑdqó :Nòx=ö>r   ry   r   )$Útypingr   r   r   r   r   rZ  Útriton.language.semanticr   Ú r
   r|   Ú_layoutsr   r   r   r   r   r   r   Útriton._C.libtriton.gluon_irr   r   Útriton.compiler.code_generatorr   r   r   rŽ   ré  rS   r   r,   re   Ú__triton_builtin__rg   ry   r   r   r   ú<module>ró     sŽ   ðß ;Õ ;Û Ý 3Ý ÷ M÷  Mñ  Mß Pß Tá�:Ó€ð <Fñ !�ð !˜x¨¨C¨Ñ0ó !ò
Ró:ðz /3Ð Ô +÷	Mñ 	Mô~>�N 8Ñ,õ ~>r   