ó
    “~i�n  ã                   ó   • S SK JrJrJrJrJr  S SKrS SKJr  SSK	J
r  SSKJrJrJrJrJrJr  S SKJrJr  S SKJrJr  \" S	5      r\4S
\S\/ \4   4S jjrS rSS jrS\l          " S S5      r! " S S\\   5      r"g)é    )ÚSequenceÚListÚTypeVarÚTupleÚCallableN)ÚTritonSemanticé   )Ú_core)Ú
AutoLayoutÚDistributedLayoutÚDistributedLinearLayoutÚSliceLayoutÚSharedLayoutÚCoalescedLayout)ÚGluonOpBuilderÚcompute_tmem_reg_layout)Úflatten_values_to_irÚunflatten_ir_valuesÚTensorTyÚcondÚmsg_fnc                 ó,   • U (       d  U" U" 5       5      eg ©N© )r   r   Úcategorys      Úi/home/mande/repo/quber/.venv/lib/python3.13/site-packages/triton/experimental/gluon/language/_semantic.pyÚ_checkr      s   € ÞÙ‘v“xÓ Ð ð ó    c                 óT   • [        U [        5      =(       a    [        S U  5       5      $ )Nc              3   óB   #   • U  H  n[        U[        5      v •  M     g 7fr   ©Ú
isinstanceÚint)Ú.0Úis     r   Ú	<genexpr>Ú_is_int_list.<locals>.<genexpr>   s   é € Ð.QÊ5Àa¬z¸!¼S×/AÐ/AÊ5ùó   ‚)r"   r   Úall©Úvalues    r   Ú_is_int_listr,      s   € Ü�eœXÓ&×Q¬3Ñ.QÉ5Ó.QÓ+QÐQr   c                 óŒ  ^^^^^^• [        [        T[        5      S 5        [        TS;   U4S j5        [        [        T[        5      U4S j5        [        TS:¬  =(       a    TTS-
  -  S:H  S 5        [	        T5      m[        [        S	 T 5       5      U4S
 j5        [        T5      n[        US:H  S 5        Uc  / nTS:H  nU(       a  SOTmU(       a"  U H  n[        [        U5      U:H  S 5        M     [        U TUTTU5      m[        TS LUUU4S j5        U(       Ga  TS   n	TR                  (       dX  [        TR                  S   SU	S-  /:H  U4S j5        TR                  R                  SU	S-  /5        SS/TR                  S'   T$ TR                  S   SU	S-  /:w  a™  U R                  mS[        TR                  5      -  n
[        U
ST-  :„  U4S j5        TR                  nS HC  n[        TU5      n[        U5       H%  u  pèUSU	S-  /:X  d  M  XÞ   US   sUS'   XÞ'   Ts  s  $    ME      SU 35       eT$ )Nc                  ó   • g)Nzinstr_variant must be a stringr   r   r   r   Ú<lambda>Ú*_compute_tmem_reg_layout.<locals>.<lambda>   s   € Ð3Sr   )Ú32x32bÚ16x64bÚ16x128bÚ16x256bÚ16x32bx2Ú32x32b_splitnc                  ó   >• ST  3$ )Nzunknown instr_variant: r   )Úinstr_variants   €r   r/   r0      s   ø€ Ð,¨]¨OÑ<r   c                  ó"   >• S[        T 5      < 3$ )Nz!num_warps must be an int but got ©Útype©Ú	num_warpss   €r   r/   r0      s   ø€ Ð1RÔSWÐXaÓSbÑReÑ/fr   é   r	   r   c                  ó   • g)Nz)num_warps must be a power of two and >= 4r   r   r   r   r/   r0      s   € ÐJur   c              3   óB   #   • U  H  n[        U[        5      v •  M     g 7fr   r!   )r$   Údims     r   r&   Ú+_compute_tmem_reg_layout.<locals>.<genexpr>   s   é € Ð5ªu¨Œz˜#œs×#Ð#ªuùr(   c                  ó   >• ST  3$ )Nz#shape entries must be ints but got r   ©Úshapes   €r   r/   r0      s   ø€ ÐAdÐejÐdkÑ?lr   é   c                  ó   • g)Nzexpected a 2D tensorr   r   r   r   r/   r0      s   € Ð4r   r6   r1   c                  ó   • g)Nzcga_layout basis rank mismatchr   r   r   r   r/   r0   (   s   € Ð/Or   c                  ó   >• ST  ST ST 3$ )NzTMEM layout 'z' unsupported for shape z and num_warps r   )Úatom_variantr=   rE   s   €€€r   r/   r0   3   s   ø€ �] < .Ð0HÈÈÈÐ_hÐ^iÑjr   éÿÿÿÿc                  ó   >• ST  3$ )NzJsplitn with 1 register requires the last lane basis to be [0, N / 2]. Got r   )Ú
layout_objs   €r   r/   r0   ;   s   ø€ ÐgÐhrÐgsÑtr   é    c                  ó   >• SST -   ST  S3$ )NzETo be able to `tmem.load` into `tl.split` you need to have more than rN   Ú 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/   r0   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_basesrR   ÚappendÚprimitive_bitwidthÚgetattrÚ	enumerate)Ú
element_tyrE   Úlayoutr=   r8   Ú
cga_layoutÚrankÚsplitnÚbasisÚNÚnum_regrW   Ú	bases_strÚbasesr%   rJ   rQ   rM   s    ` ``          @@@r   Ú_compute_tmem_reg_layoutrf      sS  ý€ Ü
Œ:�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æÛˆEÜ”3�u“: Ñ%Ñ'OÖPñ  ô )ØØØØØØó€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Ø  A¨¡F Õ+Ø27±(¸IÀb¹MÐ/˜	 "™ u¡xØ)Ô)ó !1ñ :ð \ÐRÐSYÐRZÐ[Ó[�5ØÐr   Tc                   ó.   • \ rS rSrS\4S jrS rS rSrg)ÚGluonCallerContextéU   r=   c                 ó   • Xl         g r   r<   )Úselfr=   s     r   Ú__init__ÚGluonCallerContext.__init__W   s   € Ø"�r   c                 ó    • SU R                    3$ )NÚ_NWr<   ©rk   s    r   ÚmangleÚGluonCallerContext.mangleZ   s   € Ø�T—^‘^Ð$Ð%Ð%r   c                 óZ   • UR                  SUR                  U R                  5      5        g )Nzttg.num-warps)Úset_attrÚget_int32_attrr=   )rk   ÚfnÚbuilders      r   Úinitialize_calleeÚ$GluonCallerContext.initialize_callee]   s   € Ø
�‰�O W×%;Ñ%;¸D¿N¹NÓ%KÕLr   r<   N)	Ú__name__Ú
__module__Ú__qualname__Ú__firstlineno__r#   rl   rq   rx   Ú__static_attributes__r   r   r   rh   rh   U   s   † ð# #ô #ò&õMr   rh   c            
       ó¦  ^ • \ rS rSr% \R
                  r\r\\S'   S\4S jr	S r
S rS\\   S\\   4S jrS	\S
\S\4S jrS\S\S\4U 4S jjrS\S\\\4   4U 4S jjrS	\S\\   S\4U 4S jjrS	\S\\   S\4S jrS\S\S\4U 4S jjrU 4S jrS	\S\\   S\4U 4S jjrS rS rS@S jrS rS  rS! rS" r S# r!S$ r"S% r#S& r$S' r%S( r&S) r'S* r(S+ r)\*S, 5       r+S-\,\   S
\S.\S\\S/4   4S0 jr-S-\,\   S
\S\\S/4   4S1 jr.S	\S2\S3\S\4S4 jr/S\S\S\S\4S5 jr0S6\S7\S
\S\4S8 jr1S6\S\4S9 jr2S:\,\   S;\,\   4S< jr3S= r4S> r5S?r6U =r7$ )AÚGluonSemanticéa   rw   c                 ó   • Xl         g r   ©rw   )rk   rw   s     r   rl   ÚGluonSemantic.__init__g   s   € Ø�r   c                 ó–   • U/ :X  a  UnO0[         R                  " X#U R                  R                  U5      5      nU R	                  X5      $ r   )ÚttglÚdistributed_typerw   Úget_gluon_layout_from_tensorÚtensor)rk   ÚhandleÚ	scalar_tyrE   Útys        r   Ú_wrap_handle_infer_layoutÚ'GluonSemantic._wrap_handle_infer_layoutj   s?   € Ø�B‹;Ø‰Bä×&Ò& y¸¿¹×9bÑ9bÐciÓ9jÓkˆBØ�{‰{˜6Ó&Ð&r   c                 óx   • U R                  UR                  UR                  R                  UR                  5      $ r   )r�   rŠ   r;   ÚscalarrE   )rk   r‰   s     r   Ú_wrap_tensor_infer_layoutÚ'GluonSemantic._wrap_tensor_infer_layoutq   s,   € Ø×-Ñ-¨f¯m©m¸V¿[¹[×=OÑ=OÐQW×Q]ÑQ]Ó^Ð^r   Ú	lhs_shapeÚ	rhs_shapec                 ó`  • [        U5      [        U5      :w  a  [        SU SU 35      e/ n[        U5       Hs  u  pEX$   nUS:X  a  UR                  U5        M"  US:X  d  Xe:X  a  UR                  U5        M@  [        S[	        U5      -   S-   [	        U5      -   S-   [	        U5      -   5      e   U$ )Nú!Cannot broadcast, rank mismatch: ú, r	   z?Cannot make_shape_compatible: incompatible dimensions at index ú: z and )rV   Ú
ValueErrorr[   rX   rT   )rk   r“   r”   Ú	ret_shaper%   ÚleftÚrights          r   Ú_broadcast_shapesÚGluonSemantic._broadcast_shapest   sË   € Üˆy‹>œS ›^Ó+ÜÐ@ÀÀÈ2ÈiÈ[ÐYÓZÐZàˆ	Ü  Ö+‰GˆAØ‘LˆEØ�q‹yØ× Ñ  Ö'Ø˜1“* %£-Ø× Ñ  Ö&ä ð "-Ü/2°1«vñ"6Ø8<ñ"=Ü?BÀ4»yñ"IØKRñ"SÜUXÐY^ÓU_ñ"`ó að añ ,ð Ðr   ÚinputÚaxisÚreturnc                 óÄ  ^^^• TR                    Vs/ s H  n[        R                  " U5      PM     nnUR                  TS5        TS:  a  T[	        TR                   5      -  m[        [        TR                  [        R                  5      U4S j5        TR                  R                  m[        [        T[        [        [        45      U4S j5        [        [        T[        [        45      =(       d    TR                  T:H  UU4S j5        U R                  R                  TR                   T5      nU R#                  UTR                  R$                  U5      $ s  snf )Nr	   r   c                  ó$   >• ST R                   < 3$ ©Nz=expected expand_dims input to be a distributed_type but got: r:   ©rŸ   s   €r   r/   Ú+GluonSemantic.expand_dims.<locals>.<lambda>Œ   ó   ø€ ÐVÐW\×WaÑWaÑVdÑer   c                  ó   >• ST  3$ )Nz;expected expand_dims input to have a SliceLayout, but got: r   ©r]   s   €r   r/   r¦   �   s   ø€ ÐTÐU[ÐT\Ñ]r   c                  ó(   >• ST  STR                    3$ )Nz7expected expand_dims input layout to be sliced in axis z	 but got ©rA   )r    r]   s   €€r   r/   r¦   ’   s   ø€ ÐMÈdÈVÐS\Ð]c×]gÑ]gÐ\hÑir   )rE   r†   Ú_unwrap_if_constexprÚinsertrV   r   r"   r;   r‡   r]   r   r   r   rA   rw   Úcreate_expand_dimsrŠ   r�   r�   )rk   rŸ   r    ÚxÚ	dst_shaperŠ   r]   s    ``   @r   Úexpand_dimsÚGluonSemantic.expand_dims„   s  ú€ Ø;@¿;º;ÓGº;°a”T×.Ò.¨qÖ1¹;ˆ	Ð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                 óž   >• U R                  X5      u  p[        UR                  / :g  S 5        [        TU ]  X5      nU R                  U5      $ )Nc                  ó   • g)NzCannot join scalars in gluonr   r   r   r   r/   Ú$GluonSemantic.join.<locals>.<lambda>™   s   € Ð&Dr   )Úbroadcast_impl_valuer   rE   ÚsuperÚjoinr‘   )rk   r³   r´   r+   Ú	__class__s       €r   rº   ÚGluonSemantic.join—   sG   ø€ Ø×(Ñ(¨Ó.‰ˆÜˆq�w‰w˜"‰}ÑDÔEÜ‘‘˜QÓ"ˆØ×-Ñ-¨eÓ4Ð4r   c                 ój   >• [         TU ]  U5      u  p#U R                  U5      U R                  U5      4$ r   )r¹   Úsplitr‘   )rk   r³   ÚlhsÚrhsr»   s       €r   r¾   ÚGluonSemantic.split�   s5   ø€ Ü‘7‘= Ó#‰ˆØ×-Ñ-¨cÓ2°D×4RÑ4RÐSVÓ4WÐWÐWr   Údimsc                 óD   >• [         TU ]  X5      nU R                  U5      $ r   )r¹   Úpermuter‘   )rk   rŸ   rÂ   r+   r»   s       €r   rÄ   ÚGluonSemantic.permute¡   s"   ø€ Ü‘‘ Ó,ˆØ×-Ñ-¨eÓ4Ð4r   rE   c                 ó   ^^^• [        [        TR                  [        R                  5      U4S j5        TR                  R                  5       m[        [        T5      [        T5      :H  UU4S j5        TT:X  a  T$ [        T5       H3  u  p4TU   U:w  d  M  US:w  d  M  [        STU    SU SU ST ST 3
5      e   [        R                  " TR                  R                  TTR                  R                  5      nU R                  R                  TR                  UR                  U R                  5      5      nU R                  Xe5      $ )	Nc                  ó$   >• ST R                   < 3$ r¤   r:   r¥   s   €r   r/   Ú4GluonSemantic.broadcast_impl_shape.<locals>.<lambda>§   r§   r   c                  ó   >• ST ST  3$ )Nr–   r—   r   )rE   Ú	src_shapes   €€r   r/   rÈ   ©   s   ø€ Ð7XÐYbÐXcÐceÐfkÐelÑ5mr   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_shapesrV   r[   r™   r�   r]   rw   Úcreate_broadcastrŠ   Úto_irr‰   )rk   rŸ   rE   r%   ÚitemÚret_tyrŠ   rÊ   s    ``    @r   Úbroadcast_impl_shapeÚ"GluonSemantic.broadcast_impl_shape¥   s"  ú€ ÜŒz˜%Ÿ*™*¤d×&;Ñ&;Ó<Üeô	gà—J‘J×/Ñ/Ó1ˆ	ÜŒs�9‹~¤ U£Ñ+Õ-mÔnØ�IÓØˆLÜ  Ö+‰GˆAØ�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Ó*Ð*r   r¿   rÀ   c                 óˆ  >^^	• UR                   mUR                   m	TR                  5       (       a  T	R                  5       (       d  [        T
U ]  X5      $ [	        [        T[        R                  5      U4S j5        [	        [        T	[        R                  5      U	4S j5        TR                  5       nT	R                  5       nU R                  X45      n[        TR                  [        5      n[        T	R                  [        5      nU(       a$  U(       d  U R                  UT	R                  5      nOjU(       a$  U(       d  U R                  UTR                  5      nO?TR                  T	R                  :w  a%  [        STR                   ST	R                   35      eU R                  X5      nU R                  X%5      nX4$ )Nc                  ó   >• ST < 3$ )Nz@expected broadcast left input to be a distributed_type but got: r   )Úlhs_tys   €r   r/   Ú4GluonSemantic.broadcast_impl_value.<locals>.<lambda>½   s   ø€ ÐYÐZ`ÑYcÑdr   c                  ó   >• ST < 3$ )NzAexpected broadcast right input to be a distributed_type but got: r   )Úrhs_tys   €r   r/   rÕ   ¿   s   ø€ ÐZÐ[aÑZdÑer   zLayout mismatch in broadcast: z vs )r;   Úis_blockr¹   r¸   r   r"   r†   r‡   rË   r�   r]   r   Úset_auto_layoutr™   rÐ   )rk   r¿   rÀ   r“   r”   rš   Úis_lhs_autoÚis_rhs_autorÔ   r×   r»   s           @@€r   r¸   Ú"GluonSemantic.broadcast_impl_valueµ   sT  ú€ Ø—‘ˆØ—‘ˆà�‰× Ñ ¨¯©×(9Ñ(9Ü‘7Ñ/°Ó9Ð9äŒz˜&¤$×"7Ñ"7Ó8Üdô	fäŒz˜&¤$×"7Ñ"7Ó8Üeô	gð ×+Ñ+Ó-ˆ	Ø×+Ñ+Ó-ˆ	Ø×*Ñ*¨9Ó@ˆ	ä  §¡´
Ó;ˆÜ  §¡´
Ó;ˆÞž{Ø×&Ñ& s¨F¯M©MÓ:‰CÞ¦Ø×&Ñ& s¨F¯M©MÓ:‰CØ�]‰]˜fŸm™mÓ+ÜÐ=¸f¿m¹m¸_ÈDÐQW×Q^ÑQ^ÐP_Ð`ÓaÐaà×'Ñ'¨Ó7ˆØ×'Ñ'¨Ó7ˆØˆxˆr   c                 óŽ   >• X!-
  /nUc
  [        5       n[        R                  " [        R                  XC5      n[        TU ]  XUS9$ )N)rÏ   )r   r†   r‡   Úint32r¹   Úarange)rk   ÚstartÚendr]   rE   rÏ   r»   s         €r   rß   ÚGluonSemantic.arangeÒ   sA   ø€ Ø‘�ˆØ‰>Ü“\ˆFÜ×&Ò&¤t§z¡z°5ÓAˆÜ‰w‰~˜e°ˆ~Ð8Ð8r   r°   Úcan_reorderc                 ój   >• [        U(       + S 5        [        TU ]	  XU5      nU R                  U5      $ )Nc                  ó   • g)Nz%can_reorder is not supported in gluonr   r   r   r   r/   Ú'GluonSemantic.reshape.<locals>.<lambda>Ú   s   € Ð(Or   )r   r¹   Úreshaper‘   )rk   rŸ   r°   rã   r+   r»   s        €r   rç   ÚGluonSemantic.reshapeÙ   s1   ø€ Ü�;ŒÑ OÔPÜ‘‘ °+Ó>ˆØ×-Ñ-¨eÓ4Ð4r   c                 ó  • [        U5      S:X  a  U$ [        R                  " UR                  X#5      nU R                  R                  UR                  U R                  5      UR                  5      n[        R                  " XT5      $ )Nr   )	rV   r†   r‡   Údtyperw   Úcreate_splatrÍ   rŠ   r‰   )rk   r+   rE   r]   rÏ   rŠ   s         r   ÚsplatÚGluonSemantic.splatÞ   s^   € Üˆu‹:˜‹?ØˆLÜ×&Ò& u§{¡{°EÓBˆØ—‘×*Ñ*¨6¯<©<¸¿¹Ó+EÀuÇ|Á|ÓTˆÜ�{Š{˜6Ó*Ð*r   c                 ób   • U R                  X#5      nUc
  [        5       nU R                  XQU5      $ r   )Úmake_scalarr   rì   )rk   rE   r+   rê   r]   r�   s         r   ÚfullÚGluonSemantic.fullå   s/   € Ø×!Ñ! %Ó/ˆØ‰>Ü“\ˆFØ�z‰z˜&¨Ó0Ð0r   c                 óâ  ^^• UR                   m[        [        T[        R                  5      U4S j5        [        [        T[        R
                  5      U4S j5        [        R                  " TR                  TR                  T5      nUR                  U R                  5      nU(       a‰  U R                  R                  XQR                  5      (       d_  [        STR                   ST SU R                  TR                  TR                  5       SU R                  TTR                  5       35      eU R                  R                  XQR                  5      n[        R                   " Xd5      $ )Nc                  ó   >• ST < 3$ )Nz@expected convert_layout input to be a distributed_type but got: r   )rŒ   s   €r   r/   Ú.GluonSemantic.convert_layout.<locals>.<lambda>î   s   ø€ ÐYÐZ\ÑY_Ñ`r   c                  ó   >• ST  3$ ©Nz4expected 'layout' to be a DistributedLayout but got r   r©   s   €r   r/   rô   ð   ó   ø€ ÐMÈfÈXÑVr   zlayout conversion from z to z) is not trivial.
The linear layouts are:
Ú
)r;   r   r"   r†   r‡   r   r\   rE   rÍ   rw   Úis_convert_layout_trivialrŠ   Ú	TypeErrorr]   Úto_linear_layoutÚcreate_convert_layoutr‰   )rk   r+   r]   Úassert_trivialrÏ   Ú	ret_ty_irrŠ   rŒ   s     `    @r   Úconvert_layoutÚ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È×T`ÑT`×"aÑ"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¿|¹|ÓLˆÜ�{Š{˜6Ó*Ð*r   c                 ó.  ^^^• [        [        T[        R                  5      U4S j5        [        [	        T5      U4S j5        [        [        T[        R
                  5      U4S j5        [        R                  " TTTT5      nUb@  U R                  R                  UR                  U R                  5      UR                  5      nO4U R                  R                  UR                  U R                  5      5      n[        R                  " UTTTT5      $ )Nc                  ó   >• ST  3$ )Nz,expected 'element_ty' to be a dtype but got r   )r\   s   €r   r/   Ú/GluonSemantic.allocate_shared.<locals>.<lambda>û   s   ø€ Ð=iÐjtÐiuÑ;vr   c                  ó   >• ST  3$ ©Nz1all elements of 'shape' must be integers but got r   rD   s   €r   r/   r  ü   ó   ø€ Ð._Ð`eÐ_fÑ,gr   c                  ó   >• ST  3$ ©Nz/expected 'layout' to be a SharedLayout but got r   r©   s   €r   r/   r  þ   ó   ø€ ÐHÈÈÑQr   )r   r"   r†   rê   r,   r   Úshared_memory_descriptor_typerw   Úcreate_local_allocrÍ   rŠ   Úshared_memory_descriptor)rk   r\   rE   r]   r+   rŒ   rŠ   s    ```   r   Úallocate_sharedÚ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                 óV  ^• [        [        T[        R                  5      U4S j5        [        R                  " UR
                  UR                  T5      nU R                  R                  UR                  U R                  5      UR                  5      n[        R                  " XC5      $ )Nc                  ó   >• ST  3$ rö   r   r©   s   €r   r/   Ú+GluonSemantic.shared_load.<locals>.<lambda>  r÷   r   )r   r"   r†   r   r‡   rê   rE   rw   Úcreate_local_loadrÍ   rŠ   r‰   )rk   Úmem_descr]   rÏ   rŠ   s     `  r   Úshared_loadÚGluonSemantic.shared_load  sp   ø€ ÜŒz˜&¤$×"8Ñ"8Ó9ÜVô	Xä×&Ò& x§~¡~°x·~±~ÀvÓNˆØ—‘×/Ñ/°·±¸T¿\¹\Ó0JÈHÏOÉOÓ\ˆÜ�{Š{˜6Ó*Ð*r   c                 óX  ^^• [        [        T[        R                  5      U4S j5        [        TR                  TR                  :H  UU4S j5        [        TR
                  TR
                  :H  UU4S j5        U R                  R                  TR                  TR                  5        g )Nc                  ó    >• S[        T 5       3$ )Nz+expected 'value' to be a tensor, but got a r:   r*   s   €r   r/   Ú,GluonSemantic.shared_store.<locals>.<lambda>  s   ø€ Ð9dÔeiÐjoÓepÐdqÑ7rr   c                  ó>   >• STR                    ST R                    S3$ )Nzsource shape z and destination shape ú must matchrD   ©r  r+   s   €€r   r/   r    ó   ø€ ˜ u§{¡{ mÐ3JÈ8Ï>É>ÐJZÐZeÑfr   c                  ó>   >• STR                    ST R                    S3$ )Nzsource dtype z and destination dtype r  ©rê   r  s   €€r   r/   r    r  r   )	r   r"   r†   r‰   rE   rê   rw   Úcreate_local_storerŠ   )rk   r  r+   s    ``r   Úshared_storeÚGluonSemantic.shared_store  sk   ù€ ÜŒz˜%¤§¡Ó-Ô/rÔsÜˆu�{‰{˜hŸn™nÑ,Ýfô	häˆu�{‰{˜hŸn™nÑ,Ýfô	hà�‰×'Ñ'¨¯©¸¿¹ÕFr   c                 ó¢  • [        U[        R                  5      (       d  [        S[	        U5       35      e[        U[        R
                  5      (       d  [        S[	        U5       35      eUR                  UR                  :w  a&  [        SUR                   SUR                   S35      eUR                  UR                  :w  a&  [        SUR                   SUR                   S35      eUR                  UR                  [        UR                  5      * S  :w  a%  [        S	UR                   S
UR                   35      eUR                  R                  U R                  5      nUR                  R                  U R                  5      nU R                  R                  X4[        UR                  5      UR                  R                   5      $ )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
  rE   r™   r\   Úalloc_shaperV   r]   Ú_to_irrw   Úget_shared_bank_conflictsrU   rY   )rk   Údistr_tyÚ	shared_tyÚreg_attrÚshared_attrs        r   Úbank_conflictsÚGluonSemantic.bank_conflicts  s§  € Ü˜(¤D×$9Ñ$9×:Ñ:ÜØ[Ô\`ÐaiÓ\jÐ[kÐlónð nô ˜)¤T×%GÑ%G×HÑ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                 ó¬  ^• [        [        T[        [        45      U4S j5        [        U[        5      (       d  [	        U5      n[
        R                  " T5      m[        T[        [        45      (       a  [
        R                  " T5      $ [
        R                  " U R                  R                  TR                  U R                  5      U5      5      $ )Nc                  ó    >• S[        T 5       3$ )Nz2Expected a DistributedLayout or SharedLayout, got r:   r©   s   €r   r/   Ú0GluonSemantic.to_linear_layout.<locals>.<lambda>1  s   ø€ ÐKÌDÐQWËLÈ>ÑZr   )r   r"   r   r   rU   r†   r¬   r   r   Ú	constexprrw   rû   r$  )rk   r]   rE   s    ` r   rû   ÚGluonSemantic.to_linear_layout/  s–   ø€ ÜŒz˜&Ô#4´lÐ"CÓDÜZô	\ô ˜%¤×&Ñ&Ü˜“KˆEä×*Ò*¨6Ó2ˆä�fœzÔ+BÐC×DÑDÜ—>’> &Ó)Ð)ä�~Š~˜dŸl™l×;Ñ;¸F¿M¹MÈ$Ï,É,Ó<WÐY^Ó_Ó`Ð`r   c                 óN   • U R                   R                  UR                  5        g r   )rw   Úcreate_local_deallocrŠ   )rk   r  s     r   Úshared_deallocÚGluonSemantic.shared_dealloc=  s   € Ø�‰×)Ñ)¨(¯/©/Õ:r   c                 ó¢  ^^• TR                   n[        [        T[        5      U4S j5        [        [        UR                  [
        5      U4S j5        U R                  R                  TR                  U R                  5      TR                  5      n[        R                  " UR                  UR                  T5      nU R                  XE5      $ )Nc                  ó   >• ST  3$ )Nz9set_auto_layout must set to a distributed layout but got r   r©   s   €r   r/   Ú/GluonSemantic.set_auto_layout.<locals>.<lambda>C  s   ø€ ÐRÐSYÐRZÑ[r   c                  ó6   >• ST R                   R                   3$ )Nz4set_auto_layout input must have auto layout but got )r;   r]   r*   s   €r   r/   r7  E  s   ø€ ÐMÈeÏjÉj×N_ÑN_ÐM`Ñar   )r;   r   r"   r   r]   r   rw   Úcreate_set_auto_layoutr$  rŠ   r†   r‡   r\   rE   r‰   )rk   r+   r]   Úsrc_tyrŠ   Úres_tys    ``   r   rÙ   ÚGluonSemantic.set_auto_layout@  s’   ù€ Ø—‘ˆÜŒz˜&Ô"3Ó4Ü[ô	]äŒz˜&Ÿ-™-¬Ó4Üaô	cà—‘×4Ñ4°V·]±]À4Ç<Á<Ó5PÐRW×R^ÑR^Ó_ˆÜ×&Ò& v×'8Ñ'8¸&¿,¹,ÈÓOˆØ�{‰{˜6Ó*Ð*r   c                 óP  ^^^• [        [        T[        5      U4S j5        [        [        T[        5      U4S j5        [        [        T[        5      U4S j5        S/UR                  -  nTUT'   [	        UR
                  5      nTUT'   UR                  n[        R                  " UR                  XgUR                  R                  5      nU R                  n	U	R                  UR                  U	5      UR                  U5      n
[        R                   " U
40 UR"                  D6$ )Nc                  ó   >• ST  3$ )Nz&expected 'start' to be an int but got r   )rà   s   €r   r/   Ú-GluonSemantic.memdesc_slice.<locals>.<lambda>K  s   ø€ Ð1WÐX]ÐW^Ñ/_r   c                  ó   >• ST  3$ )Nz'expected 'length' to be an int but got r   )Úlengths   €r   r/   r?  L  s   ø€ Ð2YÐZ`ÐYaÑ0br   c                  ó   >• ST  3$ )Nz$expected 'dim' to be an int but got r   r«   s   €r   r/   r?  M  s   ø€ Ð/SÐTWÐSXÑ-Yr   r   )r   r"   r#   r_   rU   rE   r]   r†   r
  rê   r;   r#  rw   Úcreate_memdesc_subslicerÍ   rŠ   r  Ú__dict__)rk   r  rà   rA  rA   ÚoffsetsrE   r]   rŒ   rw   rŠ   s     ```      r   Úmemdesc_sliceÚGluonSemantic.memdesc_sliceJ  sÙ   ú€ ÜŒz˜%¤Ó%Ô'_Ô`ÜŒz˜&¤#Ó&Ô(bÔcÜŒz˜#œsÓ#Ô%YÔZØ�#˜Ÿ™Ñ%ˆØˆ�‰Ü�X—^‘^Ó$ˆØˆˆc‰
Ø—‘ˆÜ×/Ò/°·±ÀÈxÏ}É}×OhÑOhÓiˆØ—,‘,ˆØ×0Ñ0°·±¸'Ó1BÀHÇOÁOÐU\Ó]ˆÜ×,Ò,¨VÑC°r·{±{ÑCÐCr   c                 óÞ  ^• U R                  T5      m[        TR                  [        R                  :H  U4S j5        UR
                  SS  nU R                  T5      R                  mUR                  n[        R                  " UR                  X4U5      nU R                  nUR                  UR                  U5      UR                  T5      n[        R                  " U40 UR                  D6$ )Nc                  ó"   >• ST R                    3$ )Nz%expected 'index' to be int32 but got r:   ©Úindexs   €r   r/   Ú-GluonSemantic.memdesc_index.<locals>.<lambda>Z  s   ø€ Ð3XÐY^×YcÑYcÐXdÑ1er   r	   )Ú	to_tensorr   r;   r†   rÞ   rE   rŠ   r]   r
  rê   rw   Úcreate_memdesc_indexrÍ   r  rD  )rk   r  rK  rE   r]   rŒ   rw   rŠ   s     `     r   Úmemdesc_indexÚGluonSemantic.memdesc_indexX  s³   ø€ Ø—‘˜uÓ%ˆÜˆu�z‰zœTŸZ™ZÑ'Ô)eÔfØ—‘˜q˜rÐ"ˆØ—‘˜uÓ%×,Ñ,ˆØ—‘ˆÜ×/Ò/°·±ÀÈuÓUˆØ—,‘,ˆØ×-Ñ-¨b¯h©h°wÓ.?ÀÇÁÐRWÓXˆÜ×,Ò,¨VÑC°r·{±{ÑCÐCr   c                 ór  ^^• [        [        T5      U4S j5        [        [        T5      [        TR                  5      :H  UU4S j5        T Vs/ s H  nTR                  U   PM     nnTR                  R
                  nUS [        U5      TR                  -
   nUT Vs/ s H   o5[        U5      TR                  -
  S  U   PM"     sn-  nU R                  R                  TR                  T5      nU R                  R                  U5      n[        R                  " UTR                  UXhS9$ s  snf s  snf )Nc                  ó   >• ST  3$ )Nz1all elements of 'order' must be integers but got r   )Úorders   €r   r/   Ú-GluonSemantic.memdesc_trans.<locals>.<lambda>d  r  r   c                  ó<   >• ST R                    S[        T5       S3$ )Nzsource rank (z) and order length (z) must match)r_   rV   )r  rS  s   €€r   r/   rT  g  s   ø€ �m H§M¡M ?Ð2FÄsÈ5ÃzÀlÐR^Ñ_r   ©r\   rE   r#  r]   )r   r,   rV   rE   r;   r#  r_   rw   Úcreate_memdesc_transrŠ   Úget_gluon_layout_from_memdescr†   r  rê   )	rk   r  rS  r%   rE   r#  Únew_alloc_shaperŠ   r]   s	    ``      r   Úmemdesc_transÚGluonSemantic.memdesc_transc  s  ù€ ÜŒ|˜EÓ"Ô$gÔhÜÜ�‹Jœ#˜hŸn™nÓ-Ñ-Ý_ô	añ -2Ó2ªE q�—‘ Ô"©EˆÐ2Ø—m‘m×/Ñ/ˆØ%Ð&G¤s¨;Ó'7¸(¿-¹-Ñ'GÐHˆØÑW\Ó]ÒW\ÐRS¬¨KÓ(8¸8¿=¹=Ñ(HÐ(IÐJÈ1ÔMÑW\Ñ]Ñ]ˆà—‘×2Ñ2°8·?±?ÀEÓJˆØ—‘×;Ñ;¸FÓCˆÜ×,Ò,¨VÀÇÁÐV[Ø9HñYð 	Yùò 3ùò ^s   ÁD/Â!'D4c                 ó  ^^• [        [        T5      U4S j5        [        [        R                  " T5      [        R                  " TR                  5      :H  UU4S j5        U R
                  R                  TR                  T5      nU R
                  R                  U5      nTR                  R                  n[        U5      TR                  -
  nUS U [        T5      -   n[        R                  " UTR                   TUUS9$ )Nc                  ó   >• ST  3$ r  r   rD   s   €r   r/   Ú/GluonSemantic.memdesc_reshape.<locals>.<lambda>t  r  r   c                  ó(   >• ST R                    ST 3$ )Nz)memdesc_reshape total elements mismatch: z -> rD   )r  rE   s   €€r   r/   r^  w  s   ø€ Ð@Ø Ÿ™Ð' t¨E¨7ñ4r   rV  )r   r,   ÚmathÚprodrE   rw   Úcreate_memdesc_reshaperŠ   rX  r;   r#  rV   r_   rU   r†   r  rê   )rk   r  rE   rŠ   r]   r#  Ú
prefix_lenrY  s    ``     r   Úmemdesc_reshapeÚGluonSemantic.memdesc_reshapes  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[        R                  5      U4S j5        [        [	        T5      U4S j5        [        [        T[        R
                  5      U4S j5        [        R                  " TTTT5      nU R                  R                  UR                  U R                  5      UR                  5      n[        R                  " U40 UR                  D6$ )Nc                  ó   >• ST  3$ )Nz'expected 'dtype' to be a dtype but got r   r  s   €r   r/   Ú3GluonSemantic.memdesc_reinterpret.<locals>.<lambda>Š  s   ø€ Ð8_Ð`eÐ_fÑ6gr   c                  ó   >• ST  3$ r  r   rD   s   €r   r/   rh  ‹  r  r   c                  ó   >• ST  3$ r  r   r©   s   €r   r/   rh  �  r	  r   )r   r"   r†   rê   r,   r   r
  rw   Úcreate_memdesc_reinterpretrÍ   rŠ   r  rD  )rk   r  rê   rE   r]   rŒ   rŠ   s     ```  r   Úmemdesc_reinterpretÚ!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                 óf   • U(       a  [         R                  " X#U5      nOUnU R                  X5      $ r   )r†   r‡   r‰   )rk   r¯   r‹   rš   r]   r;  s         r   Úwrap_tensorÚGluonSemantic.wrap_tensor’  s+   € ÞÜ×*Ò*¨9ÀÓH‰FàˆFØ�{‰{˜1Ó%Ð%r   c                 ó,  ^^^• U  H5  m[        [        TR                  [        R                  5      U4S j5        M7     U  Vs/ s H  oR                  R
                  PM     snmTS   m[        [        U4S jTSS   5       5      U4S j5        g s  snf )Nc                  ó$   >• ST R                   < 3$ ©Nz#expected distributed_type but got: r:   )r¯   s   €r   r/   Ú2GluonSemantic._check_same_layout.<locals>.<lambda>œ  s   ø€ ÐHkÐlm×lrÑlrÑkuÑFvr   r   c              3   ó,   >#   • U  H	  oT:H  v •  M     g 7fr   r   )r$   ÚlÚl0s     €r   r&   Ú3GluonSemantic._check_same_layout.<locals>.<genexpr>Ÿ  s   øé € Ð0¢K˜q˜–7¢Kùs   ƒr	   c                  ó   >• ST  3$ )Nz3Expected inputs to have matching layouts, but got: r   )Úlayoutss   €r   r/   rt     s   ø€ ÐLÈWÈIÑVr   )r   r"   r;   r†   r‡   r]   r)   )Úxsr¯   rw  rz  s    `@@r   Ú_check_same_layoutÚ GluonSemantic._check_same_layout™  sp   ú€ ãˆAÜ”:˜aŸf™f¤d×&;Ñ&;Ó<Ô>vÖwñ á*,Ó-ª" Q—6‘6—=”=©"Ñ-ˆØ�Q‰ZˆÜŒsÔ0 G¨A¨B¡KÓ0Ó0ÜVõ	Xùò .s   ÁBÚinputsÚreverse.c                 óþ  ^ ^^^• TS   R                   R                  m[        T5      nU* Us=::  a  U:  d  O   SU SU S35       eUS:  a  X%-  nT H$  nUR                   R                  T:X  a  M   S5       e   T R                  R	                  T Vs/ s H  ofR
                  PM     snX$5      mU" T5        TR                  5       (       d   e[        UUU U4S j[        [        T5      5       5       5      $ s  snf )Nr   z
scan axis z must be < inputs rank (Ú)z(all scan inputs must have the same shapec              3   ó–   >#   • U  H>  nTR                  TR                  U5      TU   R                  R                  T5      v •  M@     g 7fr   ©r�   Ú
get_resultr;   r�   )r$   r%   r~  Úscan_oprk   rE   s     €€€€r   r&   Ú1GluonSemantic.associative_scan.<locals>.<genexpr>³  sD   øé € ð )â'�ð ×*Ñ*¨7×+=Ñ+=¸aÓ+@À&ÈÁ)Ç.Á.×BWÑBWÐY^×_Ð_Ú'ùó   ƒAA	)	r;   rE   rV   rw   Úcreate_scanrŠ   ÚverifyÚtupleÚrange)	rk   r~  r    Úregion_builder_fnr  r_   Útr…  rE   s	   ``     @@r   Úassociative_scanÚGluonSemantic.associative_scan¢  sì   û€ à�q‘	—‘×$Ñ$ˆÜ�5‹zˆàˆu˜Õ#˜tÕ#ÐW z°$°Ð7OÐPTÈvÐUVÐ%WÓWÐ#à�!‹8Ø‰LˆDãˆAØ—6‘6—<‘< 5Õ(ÐTÐ*TÓTÐ(ñ ð —,‘,×*Ñ*¹fÓ+Eºf¸¯H¬H¹fÑ+EÀtÓUˆÙ˜'Ô"Ø�~‰~×ÑÐÐä÷ )äœ3˜v›;Ô'ó)ó )ð 	)ùò	 ,Fs   ÂC:c                 óœ  ^ ^^^^^	^
• Tc  [        U 4S jT 5       5      mSmTS   R                  R                  m
[        T
5      m[	        STs=:*  =(       a    T:  Os  UU4S j5        T R                  T5        [        T
5       VVs/ s H  u  pEUT:w  d  M  UPM     snnm	[        U
4S jT 5       5      (       d   S5       eT R                  R                  T Vs/ s H  ofR                  PM     snT5      mU" T5        TR                  5       (       d   e[        UUU	U 4S j[        [        T5      5       5       5      $ s  snnf s  snf )Nc              3   ón   >#   • U  H*  nTR                  XR                  R                  /S S9v •  M,     g7f)F)rã   N)rç   Únumelr+   )r$   r�  rk   s     €r   r&   Ú*GluonSemantic.reduction.<locals>.<genexpr>¹  s+   øé € Ð_ÒX^ÐST˜4Ÿ<™<¨¯G©G¯M©M¨?È˜<ÕNÒX^ùs   ƒ25r   c                  ó   >• ST ST  3$ )Nz/expected reduction axis to be in the range [0, z
) but got r   ©r    r_   s   €€r   r/   Ú)GluonSemantic.reduction.<locals>.<lambda>¾  s   ø€ Ð+ZÐ[_ÐZ`Ð`jÐkoÐjpÑ)qr   c              3   óT   >#   • U  H  oR                   R                  T:H  v •  M     g 7fr   )r;   rE   )r$   r�  rE   s     €r   r&   r“  Á  s   øé € Ð9²&¨Q—6‘6—<‘< 5Ö(²&ùs   ƒ%(z-all reduction inputs must have the same shapec              3   ó–   >#   • U  H>  nTR                  TR                  U5      TU   R                  R                  T5      v •  M@     g 7fr   rƒ  )r$   r%   r~  Ú	reduce_oprš   rk   s     €€€€r   r&   r“  Ç  sD   øé € ð )â'�ð ×*Ñ*¨9×+?Ñ+?ÀÓ+BÀFÈ1ÁIÇNÁN×DYÑDYÐ[d×eÐeÚ'ùr‡  )rŠ  r;   rE   rV   r   r|  r[   r)   rw   Úcreate_reducerŠ   r‰  r‹  )rk   r~  r    rŒ  r%   Úsr�  r_   r™  rš   rE   s   ```    @@@@r   Ú	reductionÚGluonSemantic.reduction·  s  þ€ Ø‰<ÜÔ_ÑX^Ó_Ó_ˆFØˆDà�q‘	—‘×$Ñ$ˆÜ�5‹zˆÜˆq�D×Ó˜4ÔÕ!qÔrØ×Ñ Ô'Ü#,¨UÔ#3ÔAÒ#3™4˜1°q¸D±y—QÑ#3ÒAˆ	ÜÔ9±&Ó9×9Ñ9ÐjÐ;jÓjÐ9à—L‘L×.Ñ.Á&Ó/IÂ&¸Q·´Á&Ñ/IÈ4ÓPˆ	Ù˜)Ô$Ø×Ñ×!Ñ!Ð!Ð!ä÷ )äœ3˜v›;Ô'ó)ó )ð 	)ùó Bùò 0Js   Â	EÂEÃE	Únum_binsÚmaskc                 ó  • [        [        UR                  5      S:H  S 5        [        UR                  R	                  5       S 5        [        US LS 5        UbN  U R                  X15      u  p1[        UR                  R                  R                  5       S 5        UR                  nUR                  U R                  5      nU R                  R                  UR                  X#U5      nU R                  U[        R                  U/U5      $ )Nr	   c                  ó   • g)Nz histogram only supports 1D inputr   r   r   r   r/   Ú)GluonSemantic.histogram.<locals>.<lambda>Ì  s   € Ð.Pr   c                  ó   • g)Nz%histogram only supports integer inputr   r   r   r   r/   r¢  Í  s   € Ð-Tr   c                  ó   • g)Nz'histogram requires a destination layoutr   r   r   r   r/   r¢  Î  s   € Ð+Tr   c                  ó   • g)Nz"Mask must have boolean scalar typer   r   r   r   r/   r¢  Ñ  s   € Ð7[r   )r   rV   rE   rê   Úis_intr¸   r;   r�   Úis_boolrŠ   r$  rw   Úcreate_histogramro  r†   rÞ   )rk   rŸ   rž  rŸ  r]   Úlayout_attrrŠ   s          r   Ú	histogramÚ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È[ÓYˆØ×Ñ ¬¯
©
°X°JÀÓGÐGr   c           	      óÈ  • [        US LS 5        [        US 5        [        [        UR                  5      S:H  S 5        [        R                  " UR
                  R                  UR                  S   UR                  S   -   /U5      nU R                  U R                  R                  UR                  UR                  UR                  U R                  5      5      U5      $ )Nc                  ó   • g)Nz!cat requires a destination layoutr   r   r   r   r/   Ú#GluonSemantic.cat.<locals>.<lambda>Ø  s   € Ð+Nr   c                  ó   • g)Nz;current implementation of `cat` always may reorder elementsr   r   r   r   r/   r®  Ù  s   € Ð$ar   r	   c                  ó   • g)Nzcat requires a rank-1 inputr   r   r   r   r/   r®  Ú  s   € Ð,Ir   r   )r   rV   rE   r†   r‡   r;   r�   r‰   rw   Ú
create_catrŠ   rÍ   )rk   r¿   rÀ   rã   r]   Úret_types         r   ÚcatÚ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   ÚsrcrK  c                 óæ  ^^^^• [        [        TR                  [        R                  5      U4S j5        [        [        TR                  [        R                  5      U4S j5        [        TR                  R
                  R                  5       U4S j5        [        TR                  R                  5      m[        [        TR                  R                  5      T:H  S 5        [        T* Ts=:*  =(       a    T:  Os  UU4S j5        TS:  a  TT-  m[        T5       HL  nUT:X  a  M  [        TR                  R                  U   TR                  R                  U   :H  U4S j5        MN     U R                  R                  TR                  TR                  T5      nU R                  UTR                  R
                  TR                  R                  TR                  R                  5      $ )Nc                  ó$   >• ST R                   < 3$ rs  r:   )rµ  s   €r   r/   Ú&GluonSemantic.gather.<locals>.<lambda>ß  s   ø€ ÐFiÐjm×jrÑjrÑiuÑDvr   c                  ó$   >• ST R                   < 3$ rs  r:   rJ  s   €r   r/   r¸  á  s   ø€ Ð<¸U¿Z¹Z¹NÑKr   c                  ó8   >• ST R                   R                  < 3$ )Nz&expected integer scalar type but got: )r;   r�   rJ  s   €r   r/   r¸  â  s   ø€ Ð5[Ð\a×\fÑ\f×\mÑ\mÑ[pÑ3qr   c                  ó   • g)Nz0source and index tensors must have the same rankr   r   r   r   r/   r¸  å  s   € Ð6hr   c                  ó   >• ST  ST S3$ )Nzgather axis z must be < source rank (r�  r   r•  s   €€r   r/   r¸  æ  s   ø€ ¨|¸D¸6ÐAYÐZ^ÐY_Ð_`Ñ-ar   r   c                  ó   >• ST  S3$ )Nz
index dim z( must match the corresponding source dimr   )r    s   €r   r/   r¸  ï  s   ø€ ˜* T FÐ*RÑSr   )r   r"   r;   r†   r‡   r�   r¦  rV   rE   r‹  rw   Úcreate_gatherrŠ   ro  r]   )rk   rµ  rK  r    ÚdÚgatherr_   s    ```  @r   rÀ  ÚGluonSemantic.gatherÞ  sS  û€ ÜŒ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                 óü   • U R                   R                  UR                  UR                  U R                   5      U5      n[	        UR
                  R                  5      nXS==   S-  ss'   U R                  XBU5      $ )NrF   )rw   Úcreate_fp4_to_fprŠ   rÍ   rU   r;   rE   r�   )rk   rµ  Ú	elem_typer    ÚresultrE   s         r   Ú	fp4_to_fpÚGluonSemantic.fp4_to_fpô  s]   € Ø—‘×.Ñ.¨s¯z©z¸9¿?¹?È4Ï<É<Ó;XÐZ^Ó_ˆÜ�S—X‘X—^‘^Ó$ˆØ‹�qÑ‹Ø×-Ñ-¨fÀÓGÐGr   Úworker_num_warpsÚworker_num_regsc           	      ó.  ^• U H4  u  nm[        [        T[        [        R                  45      U4S j5        M6     [	        U5      S:¼  d   S5       eUS   u  pg[	        U5      S-
  nUSS  n	U[	        U5      :X  d   SU S[	        U5       S35       eU[	        U5      :X  d   SU S[	        U5       S35       eU R
                  n
U
R                  5       nU
R                  5       nU
R                  U5        UR                  Xg0 S	9n/ nUb  [        U5      nU
R                  U5        U Vs/ s H  oÿR                  5       PM     nnU	 VVs/ s H  u  nn[        U5      PM     nnn[        U/ 5      nU
R                  U5        U
R                  UUU5      nUR!                  5       R#                  U5        UR%                  U5        U
R'                  UR)                  5       / 5        U
R+                  U5      nU Vs/ s H  nUR                  5       PM     nnSn[-        U	5       HÈ  u  nu  nm[/        UU   S
9nU
R'                  UR1                  U5      U5      nUU   n[3        [	        U5      5       Vs/ s H  nUR5                  UU-   5      PM     nn[7        UT Vs/ s H  nUR8                  PM     sn5      nUR                  UU0 US9  U
R;                  5         U[	        U5      -  nMÊ     U
R=                  UR?                  5       5        [3        [	        U5      5       Vs/ s H  nURA                  U5      PM     nnUc  g [        [7        Xí Vs/ s H  oÿR8                  PM     sn5      5      $ s  snf s  snnf s  snf s  snf s  snf s  snf s  snf )Nc                  ó    >• S[        T 5       3$ )Nz9function arguments must be a tuple of arguments, but got r:   )Úargss   €r   r/   Ú/GluonSemantic.warp_specialize.<locals>.<lambda>þ  s   ø€ ÐVÔW[Ð\`ÓWaÐVbÑcr   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†   rV   rw   Ú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_partitionsr[   rh   Ú
get_regionr‹  Úget_argumentr   r;   Úcreate_warp_returnÚset_insertion_point_afterÚget_operationr„  )rk   Úfunctions_and_argsrÈ  rÉ  Ú	generatorÚ_Údefault_partitionÚdefault_argsÚnum_partitionsÚworkersrw   Ú	insert_ptÚdefault_blockÚdefault_resultsÚ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_specializeÚGluonSemantic.warp_specializeú  s   ø€ ã)‰GˆAˆtÜ”:˜d¤U¬D¯J©JÐ$7Ó8Ücöeñ *ô Ð%Ó&¨!Ó+ÐgÐ-gÓgÐ+Ø*<¸QÑ*?Ñ'ÐÜÐ/Ó0°1Ñ4ˆØ$ Q RÐ(ˆà¤Øó"
ó 
ð 	fà! .Ð!1Ð1AÄ#ÐFVÓBWÐAXÐXdÐeó	fð 
ð ¤Øó"
ó 
ð 	ià! .Ð!1Ð1AÄ#ÀoÓBVÐAWÐWgÐhó	ið 
ð —,‘,ˆØ×/Ñ/Ó1ˆ	ð  ×)Ñ)Ó+ˆØ×,Ñ,¨]Ô;Ø#×4Ñ4Ð5FÐ]_Ð4Ð`ˆØˆØÑ&Ü/°Ó@ˆLØ×!Ñ! ,Ô/Ù.:Ó;ªl¨Ÿ
™
ž©lˆÐ;ñ BIÔIÂ±g°a¸Ô+¨DÖ1ÁˆÑIÜ˜ RÓ(ˆ	Ø×'Ñ'¨	Ô2Ø×.Ñ.¨|¸YÐHXÓYˆØ× Ñ Ó"×,Ñ,¨]Ô;Ø×%Ñ% oÔ6ð 	×(Ñ(¨×)FÑ)FÓ)HÈ"ÔMØ×AÑAÀ.ÓQˆÙ/8Ó9ªy¨�S—\‘\–^©yˆ	Ð9ØˆÜ(¨Ö1‰OˆA‰|��dÜ/Ð:JÈ1Ñ:MÑNˆNØ×4Ñ4°]×5MÑ5MÈaÓ5PÐR[Ó\ˆEØ# A™ˆIÜBGÌÈIËÔBWÓXÒBW¸Q˜%×,Ñ,¨V°a©ZÖ8ÑBWˆJÐXÜ,¨ZÉdÓ9SÊdÀs¸#¿(¼(ÉdÑ9SÓTˆJØ×&Ñ& t¨ZÀÐSaÐ&ÑbØ×&Ñ&Ô(Ø”c˜)“nÑ$ŠFñ  2ð 	×)Ñ)¨%×*=Ñ*=Ó*?Ô@Ü5:¼3¸|Ó;LÔ5MÓNÒ5M°˜×(Ñ(¨Ö+Ñ5MˆÐNØÑ"ØÜÔ(¨ÈÓ7XÊÀ1¿¼ÉÑ7XÓYÓZÐZùò; <ùó Jùò :ùò YùÚ9Sùò Oùò 8Ys*   Ä*M3Å
M8Ç:M>É4NÊNÌ$NÍNc                 ój   • [         R                  " U R                  R                  R                  5      $ r   )r†   r/  rw   ÚoptionsÚnum_ctasrp   s    r   r   ÚGluonSemantic.num_ctas6  s!   € Ü�~Š~˜dŸl™l×2Ñ2×;Ñ;Ó<Ð<r   c                 ó  • UR                   bK  [        UR                   [        5      (       d   e[        R                  " UR                   R
                  5      $ [        R                  " U R                  R                  R
                  5      $ r   )rÏ  r"   rh   r†   r/  r=   rw   rÿ  )rk   rå  s     r   r=   ÚGluonSemantic.num_warps9  sb   € Ø×#Ñ#Ñ/Ü˜i×6Ñ6Ô8J×KÑKÐKÐKÜ—>’> )×":Ñ":×"DÑ"DÓEÐEÜ�~Š~˜dŸl™l×2Ñ2×<Ñ<Ó=Ð=r   rƒ   )F)8rz   r{   r|   r}   r†   r‰   Úlangr   Ú__annotations__rl   r�   r‘   r   r#   r�   r   r±   rº   r   r¾   rÄ   rÐ   r¸   rß   Úboolrç   rì   rð   rÿ   r  r  r   r*  rû   r3  rÙ   rF  rO  rZ  rd  rl  ro  Ústaticmethodr|  r   rŽ  rœ  rª  r³  rÀ  rÆ  rü  r   r=   r~   Ú__classcell__)r»   s   @r   r€   r€   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ð:[ÀHÈSÁMð :[ÐdlÐmpÑdqô :[òx=÷>ð >r   r€   r   )#Útypingr   r   r   r   r   r`  Útriton.language.semanticr   Ú r
   r†   Ú_layoutsr   r   r   r   r   r   Útriton._C.libtriton.gluon_irr   r   Útriton.compiler.code_generatorr   r   r   r™   r  rT   r   r,   rf   Ú__triton_builtin__rh   r€   r   r   r   Ú<module>r     s‚   ðß ;Õ ;Û Ý 3Ý ß x× xß Pß Tá�:Ó€ð <Fñ !�ð !˜x¨¨C¨Ñ0õ !ò
Rô:ðz /3Ð Ô +÷	Mñ 	Mô\>�N 8Ñ,õ \>r   