
    9iF                         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  S SKJr  S SKJrJr  \" S	5      r\4S
\S\/ \4   4S jjr " S S5      r " S S\\   5      rg)    )SequenceListTypeVarTupleCallableN)TritonSemantic   )_core)
AutoLayoutDistributedLayoutSliceLayout)GluonOpBuilder)flatten_values_to_irunflatten_ir_valuesTensorTycondmsg_fnc                 ,    U (       d  U" U" 5       5      eg N )r   r   categorys      l/var/www/html/land-doc-ocr/venv/lib/python3.13/site-packages/triton/experimental/gluon/language/_semantic.py_checkr      s    vx       c                   .    \ rS rSrS\4S jrS rS rSrg)GluonCallerContext   	num_warpsc                     Xl         g r   r   )selfr   s     r   __init__GluonCallerContext.__init__   s    "r   c                      SU R                    3$ )N_NWr    )r!   s    r   mangleGluonCallerContext.mangle   s    T^^$%%r   c                 Z    UR                  SUR                  U R                  5      5        g )Nzttg.num-warps)set_attrget_int32_attrr   )r!   fnbuilders      r   initialize_callee$GluonCallerContext.initialize_callee   s    
OW%;%;DNN%KLr   r    N)	__name__
__module____qualname____firstlineno__intr"   r&   r-   __static_attributes__r   r   r   r   r      s    ## #&Mr   r   c            
       H  ^  \ 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S7S 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* 5       r)S+\*\   S
\S,\S\\S-4   4S. jr+S+\*\   S
\S\\S-4   4S/ jr,S	\S0\S1\S\4S2 jr-S3\*\   S4\*\   4S5 jr.S6r/U =r0$ )8GluonSemantic   r,   c                     Xl         g r   r,   )r!   r,   s     r   r"   GluonSemantic.__init__#   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_typer,   get_gluon_layout_from_tensortensor)r!   handle	scalar_tyshapetys        r   _wrap_handle_infer_layout'GluonSemantic._wrap_handle_infer_layout&   s?    B;B&&y9b9bci9jkB{{6&&r   c                 x    U R                  UR                  UR                  R                  UR                  5      $ r   )rD   r@   typescalarrB   )r!   r?   s     r   _wrap_tensor_infer_layout'GluonSemantic._wrap_tensor_infer_layout-   s,    --fmmV[[=O=OQWQ]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 )len
ValueError	enumerateappendstr)r!   rK   rL   	ret_shapeileftrights          r   _broadcast_shapesGluonSemantic._broadcast_shapes0   s    y>S^+@2i[YZZ	 +GALEqy  '1*%-  &  "-/21v"68<"=?B4y"IKR"SUXY^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[        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: rG   r\   s   r   <lambda>+GluonSemantic.expand_dims.<locals>.<lambda>H       VW\WaWaVder   c                     > ST  3$ )Nz;expected expand_dims input to have a SliceLayout, but got: r   )layouts   r   rd   re   K   s    TU[T\]r   c                  (   > ST  STR                    3$ )Nz7expected expand_dims input layout to be sliced in axis z	 but got )dim)r]   rh   s   r   rd   re   N   s    MdVS\]c]g]g\hir   )rB   r<   _unwrap_if_constexprinsertrQ   r   
isinstancerG   r=   rh   r   r   rj   r,   create_expand_dimsr@   rD   rH   )r!   r\   r]   x	dst_shaper@   rh   s    ``   @r   expand_dimsGluonSemantic.expand_dims@   s    ;@;;G;aT..q1;	Gq!!8C$$Dz%**d&;&;<e	g""z&;
";<]	_vz*@fjjD.@i	k 00tD--fejj6G6G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      $ )NzCannot join scalars in gluon)broadcast_impl_valuer   rB   superjoinrI   )r!   rs   rt   value	__class__s       r   rx   GluonSemantic.joinS   sG    ((.qww"}<=Q"--e44r   c                 j   > [         TU ]  U5      u  p#U R                  U5      U R                  U5      4$ r   )rw   splitrI   )r!   rs   lhsrhsrz   s       r   r}   GluonSemantic.splitY   s5    7=#--c2D4R4RSV4WWWr   dimsc                 D   > [         TU ]  X5      nU R                  U5      $ r   )rw   permuterI   )r!   r\   r   ry   rz   s       r   r   GluonSemantic.permute]   s"    ,--e44r   rB   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$ ra   rb   rc   s   r   rd   4GluonSemantic.broadcast_impl_shape.<locals>.<lambda>c   rf   r   c                     > ST ST  3$ )NrN   rO   r   )rB   	src_shapes   r   rd   r   e   s    7XYbXccefkel5mr   r	   z3Cannot broadcast, the expanded size of the tensor (z ) must match the existing size (z) at non-singleton dimension rP   rO   )r   rm   rG   r<   r=   get_block_shapesrQ   rS   rR   rH   rh   r,   create_broadcastr@   to_irr?   )r!   r\   rB   rW   itemret_tyr@   r   s    ``    @r   broadcast_impl_shape"GluonSemantic.broadcast_impl_shapea   s"   z%**d&;&;<e	gJJ//1	s9~U+-mnIL +GAQx4DAI #VW\]^W_V` aCCG& I%%&Cr)Bug"? @ @ ,
 &&uzz'8'8%ARAR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   rd   4GluonSemantic.broadcast_impl_value.<locals>.<lambda>y   s    YZ`Ycdr   c                     > ST < 3$ )NzAexpected broadcast right input to be a distributed_type but got: r   )rhs_tys   r   rd   r   {   s    Z[aZder   zLayout mismatch in broadcast: z vs )rG   is_blockrw   rv   r   rm   r<   r=   r   rZ   rh   r   set_auto_layoutrR   r   )r!   r~   r   rK   rL   rV   is_lhs_autois_rhs_autor   r   rz   s           @@r   rv   "GluonSemantic.broadcast_impl_valueq   sT     (9(97/99z&$"7"78d	fz&$"7"78e	g ++-	++-	**9@	 
; 
;{&&sFMM:C&&sFMM:C]]fmm+=fmm_DQWQ^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=   int32rw   arange)r!   startendrh   rB   r   rz   s         r   r   GluonSemantic.arange   sA    >\F&&tzz5Aw~e~88r   rp   can_reorderc                 h   > [        U(       + S5        [        TU ]	  XU5      nU R                  U5      $ )Nz%can_reorder is not supported in gluon)r   rw   reshaperI   )r!   r\   rp   r   ry   rz   s        r   r   GluonSemantic.reshape   s1    ; GH+>--e44r   c                     [         R                  " UR                  X#5      nU R                  R	                  UR                  U R                  5      UR                  5      n[         R                  " XT5      $ r   )r<   r=   dtyper,   create_splatr   r@   r?   )r!   ry   rB   rh   r   r@   s         r   splatGluonSemantic.splat   sN    &&u{{EB**6<<+Eu||T{{6**r   c                 b    U R                  X#5      nUc
  [        5       nU R                  XQU5      $ r   )make_scalarr   r   )r!   rB   ry   r   rh   rH   s         r   fullGluonSemantic.full   s/    !!%/>\Fzz&00r   c                 
  ^ UR                   m[        [        T[        R                  5      U4S j5        [        R                  " TR
                  TR                  U5      nUR                  U R                  5      nU(       aF  U R                  R                  XQR                  5      (       d  [        STR                   SU S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   )rC   s   r   rd   .GluonSemantic.convert_layout.<locals>.<lambda>   s    YZ\Y_`r   zlayout conversion from z to z is not trivial)rG   r   rm   r<   r=   
element_tyrB   r   r,   is_convert_layout_trivialr@   	TypeErrorrh   create_convert_layoutr?   )r!   ry   rh   assert_trivialr   	ret_ty_irr@   rC   s          @r   convert_layoutGluonSemantic.convert_layout   s    ZZz"d334`	b&&r}}bhhGLL.	$,,"H"HT`T`"a"a5bii[VHO\]]33I||L{{6**r   c                 N   [         R                  " XX25      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                  " XaX#U5      $ r   )r<   shared_memory_descriptor_typer,   create_local_allocr   r@   shared_memory_descriptor)r!   r   rB   rh   ry   rC   r@   s          r   allocate_sharedGluonSemantic.allocate_shared   sx    //
6Q\\44RXXdll5KU\\ZF\\44RXXdll5KLF,,VPUVVr   c                    [         R                  " UR                  UR                  U5      nU R                  R                  UR                  U R                  5      UR                  5      n[         R                  " XC5      $ r   )	r<   r=   r   rB   r,   create_local_loadr   r@   r?   )r!   mem_descrh   r   r@   s        r   shared_loadGluonSemantic.shared_load   sT    &&x~~x~~vN//T\\0JHOO\{{6**r   c                 T   UR                   UR                   :X  d"   SUR                    SUR                    S35       eUR                  UR                  :X  d"   SUR                   SUR                   S35       eU R                  R                  UR                  UR                  5        g )Nzsource shape z and destination shape z must matchzsource dtype z and destination dtype )rB   r   r,   create_local_storer@   )r!   r   ry   s      r   shared_storeGluonSemantic.shared_store   s    {{hnn,}ekk]Jabjbpbpaqq|.}},{{hnn,}ekk]Jabjbpbpaqq|.}},''Fr   c                 N    U R                   R                  UR                  5        g r   )r,   create_local_deallocr@   )r!   r   s     r   shared_deallocGluonSemantic.shared_dealloc   s    ))(//:r   c                    UR                   n[        U[        5      (       d
   SU 35       e[        UR                  [        5      (       d   SUR                   R                   35       eU R
                  R                  UR                  U R
                  5      UR                  5      n[        R                  " UR                  UR                  U5      nU R                  XE5      $ )Nz9set_auto_layout must set to a distributed layout but got z4set_auto_layout input must have auto layout but got )rG   rm   r   rh   r   r,   create_set_auto_layout_to_irr@   r<   r=   r   rB   r?   )r!   ry   rh   src_tyr@   res_tys         r   r   GluonSemantic.set_auto_layout   s    &+- - 	s0ijpiq.r	s -&--$& & 	r)]^c^h^h^o^o]p'q	r &44V]]4<<5PRWR^R^_&&v'8'8&,,O{{6**r   c                    S/UR                   -  nX%U'   [        UR                  5      nX6U'   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$ )Nr   )ranklistrB   rh   r<   r   r   rG   alloc_shaper,   create_memdesc_subslicer   r@   r   __dict__)r!   r   r   lengthrj   offsetsrB   rh   rC   r,   r@   s              r   memdesc_sliceGluonSemantic.memdesc_slice   s    #%X^^$c
//x}}OhOhi,,00'1BHOOU\],,VCr{{CCr   c                    UR                   SS  nU R                  U5      R                  nUR                  n[        R
                  " UR                  X4U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$ )Nr	   )rB   	to_tensorr@   rh   r<   r   r   rG   r   r,   create_memdesc_indexr   r   r   )r!   r   indexrB   rh   rC   r,   r@   s           r   memdesc_indexGluonSemantic.memdesc_index   s    qr"u%,,//x}}OhOhi,,--bhhw.?RWX,,VCr{{CCr   c                 \   [        U5      [        UR                  5      :X  d!   SUR                   S[        U5       S35       eU Vs/ s H  o1R                  U   PM     nnUR                  R                  nUS [        U5      UR                  -
   nXb Vs/ s H   o5[        U5      UR                  -
  S  U   PM"     sn-  nU R
                  R                  UR                  U5      nU R
                  R                  U5      n[        R                  " XqR                  UXhS9$ s  snf s  snf )Nzsource rank (z) and order length (z) must matchr   rB   r   rh   )rQ   rB   r   rG   r   r,   create_memdesc_transr@   get_gluon_layout_from_memdescr<   r   r   )	r!   r   orderrW   rB   r   new_alloc_shaper@   rh   s	            r   memdesc_transGluonSemantic.memdesc_trans   s#   5zSNN  	i,X]]O;OPSTYPZ|[gh	i  -22Eq"E2mm//%&Gs;'7(--'GHW\]W\RSK(88==(H(IJ1MW\]]228??EJ;;FC,,VV[9HY 	Y 3 ^s   D$'D)c                   ^^ [        [        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 R                    ST 3$ )Nz)memdesc_reshape total elements mismatch: z -> )rB   )r   rB   s   r   rd   /GluonSemantic.memdesc_reshape.<locals>.<lambda>   s    @ 'tE74r   r   )r   mathprodrB   r,   create_memdesc_reshaper@   r   rG   r   rQ   r   r   r<   r   r   )r!   r   rB   r@   rh   r   
prefix_lenr   s    ``     r   memdesc_reshapeGluonSemantic.memdesc_reshape   s    IIe		(.. 995	
 44X__eL;;FCmm//%5
%kz2T%[@,,~~'
 	
r   c                     [         R                  " X#XC5      nU R                  R                  UR	                  U R                  5      UR
                  5      n[         R                  " U40 UR                  D6$ r   )r<   r   r,   create_memdesc_reinterpretr   r@   r   r   )r!   r   r   rB   rh   rC   r@   s          r   memdesc_reinterpret!GluonSemantic.memdesc_reinterpret  sX    //fL88$,,9OQYQ`Q`a,,VCr{{CCr   c                 f    U(       a  [         R                  " X#U5      nOUnU R                  X5      $ r   )r<   r=   r?   )r!   ro   rA   rV   rh   r   s         r   wrap_tensorGluonSemantic.wrap_tensor  s+    **9HF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: rb   )ro   s   r   rd   2GluonSemantic._check_same_layout.<locals>.<lambda>  s    HklmlrlrkuFvr   r   c              3   ,   >#    U  H	  oT:H  v   M     g 7fr   r   ).0ll0s     r   	<genexpr>3GluonSemantic._check_same_layout.<locals>.<genexpr>  s     0Kq7Ks   r	   c                     > ST  3$ )Nz3Expected inputs to have matching layouts, but got: r   )layoutss   r   rd   r     s    LWIVr   )r   rm   rG   r<   r=   rh   all)xsro   r  r  s    `@@r   _check_same_layout GluonSemantic._check_same_layout  sp    A:affd&;&;<>vw *,-"Q66=="-QZs0GABK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   rD   
get_resultrG   rH   )r   rW   r
  scan_opr!   rB   s     r   r  1GluonSemantic.associative_scan.<locals>.<genexpr>,  sD      )' **7+=+=a+@&)..BWBWY^__'   AA	)	rG   rB   rQ   r,   create_scanr@   verifytuplerange)	r!   r
  r]   region_builder_fnr  r   tr  rB   s	   ``     @@r   associative_scanGluonSemantic.associative_scan  s    q	$$5zu#t#Wz$7OPTvUV%WW#!8LDA66<<5(T*TT(  ,,**f+EfHHf+EtU'"~~ )3v;') ) 	)	 ,Fs   C:c                   ^ ^^^^^	^
 [        TS LS 5        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                      g)Nz*All-reduce is not yet implemented in gluonr   r   r   r   rd   )GluonSemantic.reduction.<locals>.<lambda>1  s    )Ur   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   rd   r  5  s    +Z[_Z``jkojp)qr   c              3   T   >#    U  H  oR                   R                  T:H  v   M     g 7fr   )rG   rB   )r   r  rB   s     r   r  *GluonSemantic.reduction.<locals>.<genexpr>8  s     9&Q66<<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   rW   r
  	reduce_oprV   r!   s     r   r  r!  >  sD      )' **9+?+?+BF1INNDYDY[dee'r  )r   rG   rB   rQ   r  rS   r  r,   create_reducer@   r  r  r  )r!   r
  r]   r  rW   sr  r   r#  rV   rB   s   ```    @@@@r   	reductionGluonSemantic.reduction0  s   t4!UVq	$$5zqD4!qr'#,U#3A#341qDyQ#3A	9&999j;jj9LL..&/I&Q&/I4P	)$!!!! )3v;') ) 	) B 0Js   >D8D8D>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   rd   )GluonSemantic.histogram.<locals>.<lambda>C  s    .Pr   c                      g)Nz%histogram only supports integer inputr   r   r   r   rd   r,  D  s    -Tr   c                      g)Nz'histogram requires a destination layoutr   r   r   r   rd   r,  E  s    +Tr   c                      g)Nz"Mask must have boolean scalar typer   r   r   r   rd   r,  H  s    7[r   )r   rQ   rB   r   is_intrv   rG   rH   is_boolr@   r   r,   create_histogramr   r<   r   )r!   r\   r(  r)  rh   layout_attrr@   s          r   	histogramGluonSemantic.histogramB  s    s5;;1$&PQu{{!!#%TUvT!#TU33D@KD499##++-/[\;;DmmDLL1..u||X[Y

XJGGr   worker_num_warpsworker_num_regsc           	         [        U5      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                  X!0 S9n/ nUb  [        U5      nU	R                  U5        U Vs/ s H  oR                  5       PM     nnU	R                  U
5        [        U5      n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[#        U5       H  n[%        UU   S9nU	R                  UR'                  U5      U5      n[#        [        U5      5       Vs/ s H  nUR)                  U5      PM     nn[+        UU Vs/ s H  nUR,                  PM     sn5      nUR                  UU   U0 US9  U	R/                  5         M     U	R1                  UR3                  5       5        [#        [        U5      5       Vs/ s H  nUR5                  U5      PM     nnUc  g [7        [+        X Vs/ s H  oR,                  PM     sn5      5      $ s  snf s  snf s  snf s  snf s  snf s  snf )Nzwarp specialize got z partitions but z warp countsz register counts)kwargsr    )r9  caller_context)rQ   r,   get_insertion_point	new_blockset_insertion_point_to_startcall_JitFunctionr   create_warp_yieldget_type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  r   
get_regionget_argumentr   rG   create_warp_returnset_insertion_point_afterget_operationr  r  )r!   default_argsdefault_partitionworker_argsworker_partitionsr6  r7  	generatornum_partitionsr,   	insert_ptdefault_blockdefault_resultsmlir_resultsrresult_types	mlir_argsws_oppartitions_oparg	arg_typesrW   r:  blockj
block_argss                             r   warp_specializeGluonSemantic.warp_specializeN  s   ./"
 
 	f!.!11A#FVBWAXXde	f 
 "
 
 	i!.!11A#oBVAWWgh	i 
 ,,//1	  ))+,,];#445F]_4`&/@L!!,/.:;l

l; 	''	2(5	..|YHXY  ",,];%%o6 	(()F)F)H"MAA.Q/89yS\\^y	9~&A/:J1:MNN44]5M5Ma5PR[\E9>s9~9NO9NA%,,Q/9NJO,Zk9Zks#((k9Z[J&&'8';ZPRcq&r&&( ' 	))%*=*=*?@5:3|;L5MN5M((+5MN"(7X17XYZZ3 < : P9Z
 O 8Ys$   
K9K&K!K&K+7K0r9   )F)1r/   r0   r1   r2   r<   r?   langr   __annotations__r"   rD   rI   r   r3   rZ   r   rq   rx   r   r}   r   r   rv   r   boolr   r   r   r   r   r   r   r   r   r   r   r   r   r   r   staticmethodr  r   r  r&  r4  rb  r4   __classcell__)rz   s   @r   r6   r6      s'   [[FD '_49 c  T T T T&5h 58 5 5Xx XE(H*<$= X5X 5U3Z 5H 5+( +5: +( +  x H :95X 5$s) 5$ 5
+
1	+W+
G
;+	DDY
*D
& X X)x'9 ) )"&)+03+?)*) 2 )# )UZ[ceh[hUi )$
Hx 
H3 
Hh 
HS[ 
H.[*23-.[JRSV-.[ .[r   r6   )typingr   r   r   r   r   r   triton.language.semanticr    r
   r<   _layoutsr   r   r   triton._C.libtriton.gluon_irr   triton.compiler.code_generatorr   r   r   rR   rf  rU   r   r   r6   r   r   r   <module>ro     sm    ; ;  3  @ @ 7 T: <F ! !xC0 !
	M 	M_[N8, _[r   