
    PiF                         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 d dlmZ d dlmZmZ  ed	          Zefd
edeg ef         fdZ G d d          Z G d dee                   ZdS )    )SequenceListTypeVarTupleCallableN)TritonSemantic   )_core)
AutoLayoutDistributedLayoutSliceLayout)GluonOpBuilder)flatten_values_to_irunflatten_ir_valuesTensorTycondmsg_fnc                 0    | s | |                      d S N )r   r   categorys      /var/www/development/aibuddy-work/election-extract/venv/lib/python3.11/site-packages/triton/experimental/gluon/language/_semantic.py_checkr      s*     !hvvxx   ! !    c                   &    e Zd ZdefdZd Zd ZdS )GluonCallerContext	num_warpsc                     || _         d S r   r   )selfr   s     r   __init__zGluonCallerContext.__init__   s    "r   c                     d| j          S )N_NWr   )r    s    r   manglezGluonCallerContext.mangle   s    %T^%%%r   c                 b    |                     d|                    | j                             d S )Nzttg.num-warps)set_attrget_int32_attrr   )r    fnbuilders      r   initialize_calleez$GluonCallerContext.initialize_callee   s,    
OW%;%;DN%K%KLLLLLr   N)__name__
__module____qualname__intr!   r$   r*   r   r   r   r   r      sR        ## # # # #& & &M M M M Mr   r   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d6d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d-f         fd.Z*d+e)e         d	ed
eed-f         fd/Z+ded0ed1ed
efd2Z,d3e)e         d4e)e         fd5Z- xZ.S )7GluonSemanticr)   c                     || _         d S r   )r)   )r    r)   s     r   r!   zGluonSemantic.__init__#   s    r   c                     |g k    r|}n.t          j        ||| j                            |                    }|                     ||          S r   )ttgldistributed_typer)   get_gluon_layout_from_tensortensor)r    handle	scalar_tyshapetys        r   _wrap_handle_infer_layoutz'GluonSemantic._wrap_handle_infer_layout&   sK    B;;BB&y%9b9bci9j9jkkB{{62&&&r   c                 X    |                      |j        |j        j        |j                  S r   )r;   r7   typescalarr9   )r    r6   s     r   _wrap_tensor_infer_layoutz'GluonSemantic._wrap_tensor_infer_layout-   s%    --fmV[=OQWQ]^^^r   	lhs_shape	rhs_shapec                    t          |          t          |          k    rt          d| d|           g }t          |          D ]\  }}||         }|dk    r|                    |           )|dk    s||k    r|                    |           Kt          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 : z and )len
ValueError	enumerateappendstr)r    r@   rA   	ret_shapeileftrights          r   _broadcast_shapeszGluonSemantic._broadcast_shapes0   s   y>>S^^++YYYiYYZZZ	 ++ 	a 	aGAtaLEqyy  ''''1**%4--  &&&&  "-/21vv"68<"=?B4yy"IKR"SUXY^U_U_"` a a ar   inputaxisreturnc                 H   d j         D             }|                    d           dk     rt          j                   z  t          t	          j        t          j                  fd           j        j        t          t	          t          t          f          fd           t          t	          t                    p
j        k    fd           | j                            j                  }|                     |j        j        |          S )Nc                 6    g | ]}t          j        |          S r   )r3   _unwrap_if_constexpr.0xs     r   
<listcomp>z-GluonSemantic.expand_dims.<locals>.<listcomp>A   s#    GGGaT.q11GGGr   r	   r   c                      d j         S Nz=expected expand_dims input to be a distributed_type but got: r=   rP   s   r   <lambda>z+GluonSemantic.expand_dims.<locals>.<lambda>H       eW\Waee r   c                      d  S )Nz;expected expand_dims input to have a SliceLayout, but got: r   )layouts   r   r^   z+GluonSemantic.expand_dims.<locals>.<lambda>K   s    ]U[]] r   c                      d  dj          S )Nz7expected expand_dims input layout to be sliced in axis z	 but got )dim)rQ   ra   s   r   r^   z+GluonSemantic.expand_dims.<locals>.<lambda>N   s    idii]c]gii r   )r9   insertrF   r   
isinstancer=   r3   r4   ra   r   r   rc   r)   create_expand_dimsr7   r;   r>   )r    rP   rQ   	dst_shaper7   ra   s    ``  @r   expand_dimszGluonSemantic.expand_dims@   s0   GG5;GGG	q!!!!88C$$$Dz%*d&;<<eeee	g 	g 	g"z&;
";<<]]]]	_ 	_ 	_vz**@fjD.@iiiii	k 	k 	k 00tDD--fej6GSSSr   abc                     |                      ||          \  }}t          |j        g k    d           t                                          ||          }|                     |          S )NzCannot join scalars in gluon)broadcast_impl_valuer   r9   superjoinr?   )r    ri   rj   value	__class__s       r   rn   zGluonSemantic.joinS   s^    ((A..1qw"}<===Q""--e444r   c                     t                                          |          \  }}|                     |          |                     |          fS r   )rm   splitr?   )r    ri   lhsrhsrp   s       r   rr   zGluonSemantic.splitY   sD    77==##S--c22D4R4RSV4W4WWWr   dimsc                 r    t                                          ||          }|                     |          S r   )rm   permuter?   )r    rP   ru   ro   rp   s       r   rw   zGluonSemantic.permute]   s.    t,,--e444r   r9   c                    t          t          j        t          j                  fd           j                                        t          t                    t                    k    fd           k    rS t                    D ];\  }}|         |k    r*|dk    r$t          d|          d| d| d d 
          <t          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>c   r_   r   c                      d d  S )NrC   rD   r   )r9   	src_shapes   r   r^   z4GluonSemantic.broadcast_impl_shape.<locals>.<lambda>e   s    5mYb5m5mfk5m5m r   r	   z3Cannot broadcast, the expanded size of the tensor (z ) must match the existing size (z) at non-singleton dimension rE   rD   )r   re   r=   r3   r4   get_block_shapesrF   rH   rG   r>   ra   r)   create_broadcastr7   to_irr6   )r    rP   r9   rL   itemret_tyr7   r{   s    ``    @r   broadcast_impl_shapez"GluonSemantic.broadcast_impl_shapea   sy   z%*d&;<<eeee	g 	g 	gJ//11	s9~~U+-m-m-m-m-mnnnIL ++ 	@ 	@GAtQx4DAII  "?W\]^W_ "? "?CG"? "?%&"? "?*3"? "?7<"? "? @ @ @ &uz'8%ARSS..u|V\\$,=W=WXX{{66***r   rs   rt   c                   	 |j         |j         	                                r	                                s"t                                          ||          S t	          t          t          j                  fd           t	          t          	t          j                  	fd                                           }	                                }| 	                    ||          }t          j
        t                    }t          	j
        t                    }|r|s|                     |	j
                  }nO|r|s|                     |j
                  }n/j
        	j
        k    rt          dj
         d	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>y   s    dZ`d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    e[aee r   zLayout mismatch in broadcast: z vs )r=   is_blockrm   rl   r   re   r3   r4   r|   rO   ra   r   set_auto_layoutrG   r   )r    rs   rt   r@   rA   rK   is_lhs_autois_rhs_autor   r   rp   s           @@r   rl   z"GluonSemantic.broadcast_impl_valueq   s      	:(9(9 	:77//S999z&$"788dddd	f 	f 	fz&$"788eeee	g 	g 	g ++--	++--	**9i@@	 
;; 
;; 	b{ 	b&&sFM::CC 	b 	b&&sFM::CC]fm++`fm``QWQ^``aaa''Y77''Y77Cxr   c                     ||z
  g}|t                      }t          j        t          j        ||          }t	                                          |||          S )N)r   )r   r3   r4   int32rm   arange)r    startendra   r9   r   rp   s         r   r   zGluonSemantic.arange   sN    u>\\F&tz5&AAww~~eS~888r   rg   can_reorderc                     t          | d           t                                          |||          }|                     |          S )Nz%can_reorder is not supported in gluon)r   rm   reshaper?   )r    rP   rg   r   ro   rp   s        r   r   zGluonSemantic.reshape   sC    ; GHHHy+>>--e444r   c                     t          j        |j        ||          }| j                            |                    | j                  |j                  }t          j        ||          S r   )r3   r4   dtyper)   create_splatr~   r7   r6   )r    ro   r9   ra   r   r7   s         r   splatzGluonSemantic.splat   sQ    &u{E6BB**6<<+E+Eu|TT{66***r   c                 |    |                      ||          }|t                      }|                     |||          S r   )make_scalarr   r   )r    r9   ro   r   ra   r>   s         r   fullzGluonSemantic.full   s;    !!%//>\\Fzz&%000r   Fc                    |j         t          t          t          j                  fd           t          j        j        j        |          }|                    | j                  }|r;| j        	                    ||j
                  st          dj         d| d          | 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    `Z\`` r   zlayout conversion from z to z is not trivial)r=   r   re   r3   r4   
element_tyr9   r~   r)   is_convert_layout_trivialr7   	TypeErrorra   create_convert_layoutr6   )r    ro   ra   assert_trivialr   	ret_ty_irr7   r:   s          @r   convert_layoutzGluonSemantic.convert_layout   s    Zz"d344````	b 	b 	b&r}bhGGLL..	 	^$,"H"HTYT`"a"a 	^\bi\\V\\\]]]33Iu|LL{66***r   c                 :   t          j        ||||          }|9| j                            |                    | j                  |j                  }n2| j                            |                    | j                            }t          j        |||||          S r   )r3   shared_memory_descriptor_typer)   create_local_allocr~   r7   shared_memory_descriptor)r    r   r9   ra   ro   r:   r7   s          r   allocate_sharedzGluonSemantic.allocate_shared   s    /
E65QQ\44RXXdl5K5KU\ZZFF\44RXXdl5K5KLLF,VZPUVVVr   c                     t          j        |j        |j        |          }| j                            |                    | j                  |j                  }t          j        ||          S r   )	r3   r4   r   r9   r)   create_local_loadr~   r7   r6   )r    mem_descra   r   r7   s        r   shared_loadzGluonSemantic.shared_load   sS    &x~x~vNN//T\0J0JHO\\{66***r   c                     |j         |j         k    sJ d|j          d|j          d            |j        |j        k    sJ d|j         d|j         d            | j                            |j        |j                   d S )Nzsource shape z and destination shape z must matchzsource dtype z and destination dtype )r9   r   r)   create_local_storer7   )r    r   ro   s      r   shared_storezGluonSemantic.shared_store   s    {hn,,,.}ek.}.}bjbp.}.}.},,,{hn,,,.}ek.}.}bjbp.}.}.},,,''FFFFFr   c                 D    | j                             |j                   d S r   )r)   create_local_deallocr7   )r    r   s     r   shared_dealloczGluonSemantic.shared_dealloc   s     ))(/:::::r   c                    |j         }t          |t                    sJ d|             t          |j        t                    sJ d|j         j                     | j                            |                    | j                  |j                  }t          j
        |j        |j        |          }|                     ||          S )Nz9set_auto_layout must set to a distributed layout but got z4set_auto_layout input must have auto layout but got )r=   re   r   ra   r   r)   create_set_auto_layout_to_irr7   r3   r4   r   r9   r6   )r    ro   ra   src_tyr7   res_tys         r   r   zGluonSemantic.set_auto_layout   s    &+- - 	s 	s.rjp.r.r	s 	s 	s&-$& & 	r 	r'q^c^h^o'q'q	r 	r 	r44V]]4<5P5PRWR^__&v'8&,OO{{66***r   c                 H   dg|j         z  }|||<   t          |j                  }|||<   |j        }t	          j        |j        |||j        j                  }| j	        }	|	
                    |                    |	          |j        |          }
t	          j        |
fi |j        S )Nr   )ranklistr9   ra   r3   r   r   r=   alloc_shaper)   create_memdesc_subslicer~   r7   r   __dict__)r    r   r   lengthrc   offsetsr9   ra   r:   r)   r7   s              r   memdesc_slicezGluonSemantic.memdesc_slice   s    #%X^$$c
/vx}Ohii,00'1B1BHOU\]],VCCr{CCCr   c                 H   |j         dd          }|                     |          j        }|j        }t	          j        |j        |||j        j                  }| j	        }|
                    |                    |          |j        |          }t	          j        |fi |j        S )Nr	   )r9   	to_tensorr7   ra   r3   r   r   r=   r   r)   create_memdesc_indexr~   r   r   )r    r   indexr9   ra   r:   r)   r7   s           r   memdesc_indexzGluonSemantic.memdesc_index   s    qrr"u%%,/vx}Ohii,--bhhw.?.?RWXX,VCCr{CCCr   c                    t          |          t          j                  k    s#J dj         dt          |           d            fd|D             }j        j        d t                    j        z
           }|fd|D             z  }| j                            j        |          }| j                            |          }t          j
        |j        |||          S )Nzsource rank (z) and order length (z) must matchc                 *    g | ]}j         |         S r   r9   )rW   rL   r   s     r   rY   z/GluonSemantic.memdesc_trans.<locals>.<listcomp>   s     222q"222r   c                 Z    g | ]'}t                    j        z
  d          |         (S r   )rF   r   )rW   rL   r   r   s     r   rY   z/GluonSemantic.memdesc_trans.<locals>.<listcomp>   s9    ]]]RSKK(8(88=(H(I(IJ1M]]]r   r   r9   r   ra   )rF   r9   r   r=   r   r)   create_memdesc_transr7   get_gluon_layout_from_memdescr3   r   r   )r    r   orderr9   new_alloc_shaper7   ra   r   s    `     @r   memdesc_transzGluonSemantic.memdesc_trans   s    5zzSN    hX]hhPSTYPZPZhhh   3222E222m/%&Gs;'7'7(-'G&GH]]]]]W\]]]]228?EJJ;;FCC,VV[9HQWY Y Y 	Yr   c                    t          t          j                  t          j        j                  k    fd           | j                            j                  }| j                            |          }j        j	        }t          |          j        z
  }|d |         t                    z   }t          j        |j        ||          S )Nc                      d j          d S )Nz)memdesc_reshape total elements mismatch: z -> r   )r   r9   s   r   r^   z/GluonSemantic.memdesc_reshape.<locals>.<lambda>   s&     4 4 4,14 4 r   r   )r   mathprodr9   r)   create_memdesc_reshaper7   r   r=   r   rF   r   r   r3   r   r   )r    r   r9   r7   ra   r   
prefix_lenr   s    ``     r   memdesc_reshapezGluonSemantic.memdesc_reshape   s    Ie	(. 9 995 5 5 5 5	
 	
 	
 44X_eLL;;FCCm/%%5
%kzk2T%[[@,~'
 
 
 	
r   c                     t          j        ||||          }| j                            |                    | j                  |j                  }t          j        |fi |j        S r   )r3   r   r)   create_memdesc_reinterpretr~   r7   r   r   )r    r   r   r9   ra   r:   r7   s          r   memdesc_reinterpretz!GluonSemantic.memdesc_reinterpret  s\    /ufeLL88$,9O9OQYQ`aa,VCCr{CCCr   c                 d    |rt          j        |||          }n|}|                     ||          S r   )r3   r4   r6   )r    rX   r8   rK   ra   r   s         r   wrap_tensorzGluonSemantic.wrap_tensor  s9     	*9iHHFFF{{1f%%%r   c                    | D ]2t          t          j        t          j                  fd           3d | D             d         t          t          fddd          D                       fd           d S )Nc                      d j         S )Nz#expected distributed_type but got: r\   )rX   s   r   r^   z2GluonSemantic._check_same_layout.<locals>.<lambda>  s    FvlmlrFvFv r   c                 &    g | ]}|j         j        S r   )r=   ra   rV   s     r   rY   z4GluonSemantic._check_same_layout.<locals>.<listcomp>  s    ---Q16=---r   r   c              3   $   K   | ]
}|k    V  d S r   r   )rW   ll0s     r   	<genexpr>z3GluonSemantic._check_same_layout.<locals>.<genexpr>  s'      00q17000000r   r	   c                      d  S )Nz3Expected inputs to have matching layouts, but got: r   )layoutss   r   r^   z2GluonSemantic._check_same_layout.<locals>.<lambda>  s    VWVV r   )r   re   r=   r3   r4   all)xsr   r   rX   s    @@@r   _check_same_layoutz GluonSemantic._check_same_layout  s     	x 	xA:afd&;<<>v>v>v>vwwww--"---QZs0000GABBK00000VVVV	X 	X 	X 	X 	Xr   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    s
J d             j                            d D             ||           |                                           sJ t           fdt          t                              D                       S )Nr   z
scan axis z must be < inputs rank ()z(all scan inputs must have the same shapec                     g | ]	}|j         
S r   r7   rW   ts     r   rY   z2GluonSemantic.associative_scan.<locals>.<listcomp>(  s    +E+E+EAH+E+E+Er   c              3      K   | ]>}                                         |          |         j        j                  V  ?d S r   r;   
get_resultr=   r>   )rW   rL   r   scan_opr    r9   s     r   r   z1GluonSemantic.associative_scan.<locals>.<genexpr>,  sa       ) ) **7+=+=a+@+@&).BWY^__) ) ) ) ) )r   )r=   r9   rF   r)   create_scanverifytuplerange)	r    r   rQ   region_builder_fnr   r   r   r   r9   s	   ``     @@r   associative_scanzGluonSemantic.associative_scan  sK   q	$5zzu####t#####%W$%W%WPT%W%W%W###!88DLD 	U 	UA6<5(((*T((((,**+E+Ef+E+E+EtWUU'"""~~ ) ) ) ) ) ) )3v;;'') ) ) ) ) 	)r   c                 b    t          d ud            d         j        j        t                    t          dcxk    ok     nc fd                                           fdt                    D             t          fdD                       s
J d             j                            d D                        |           	                                sJ t           fdt          t                              D                       S )	Nc                      dS )Nz*All-reduce is not yet implemented in gluonr   r   r   r   r^   z)GluonSemantic.reduction.<locals>.<lambda>1  s    )U r   r   c                      d d  S )Nz/expected reduction axis to be in the range [0, z
) but got r   )rQ   r   s   r   r^   z)GluonSemantic.reduction.<locals>.<lambda>5  s    )q[_)q)qko)q)q r   c                 &    g | ]\  }}|k    |S r   r   )rW   rL   srQ   s      r   rY   z+GluonSemantic.reduction.<locals>.<listcomp>7  s"    AAA41aqDyyQyyyr   c              3   8   K   | ]}|j         j        k    V  d S r   )r=   r9   )rW   r   r9   s     r   r   z*GluonSemantic.reduction.<locals>.<genexpr>8  s,      99Q16<5(999999r   z-all reduction inputs must have the same shapec                     g | ]	}|j         
S r   r   r   s     r   rY   z+GluonSemantic.reduction.<locals>.<listcomp>:  s    /I/I/IQ/I/I/Ir   c              3      K   | ]>}                                         |          |         j        j                  V  ?d S r   r   )rW   rL   r   	reduce_oprK   r    s     r   r   z*GluonSemantic.reduction.<locals>.<genexpr>>  sa       ) ) **9+?+?+B+BF1INDY[dee) ) ) ) ) )r   )r   r=   r9   rF   r   rH   r   r)   create_reducer   r   r   )r    r   rQ   r   r   r  rK   r9   s   ``` @@@@r   	reductionzGluonSemantic.reduction0  s   t4!U!UVVVq	$5zzqD4!q!q!q!q!qrrr'''AAAA9U#3#3AAA	9999&99999jj;jjjjL../I/I&/I/I/I4PP	)$$$!!!!! ) ) ) ) ) ) )3v;;'') ) ) ) ) 	)r   num_binsmaskc                 "   t          t          |j                  dk    d            t          |j                                        d            t          |d ud            |M|                     ||          \  }}t          |j        j                                        d            |j	        }|
                    | j                  }| j                            |j	        |||          }|                     |t          j        |g|          S )Nr	   c                      dS )Nz histogram only supports 1D inputr   r   r   r   r^   z)GluonSemantic.histogram.<locals>.<lambda>C  s    .P r   c                      dS )Nz%histogram only supports integer inputr   r   r   r   r^   z)GluonSemantic.histogram.<locals>.<lambda>D  s    -T r   c                      dS )Nz'histogram requires a destination layoutr   r   r   r   r^   z)GluonSemantic.histogram.<locals>.<lambda>E  s    +T r   c                      dS )Nz"Mask must have boolean scalar typer   r   r   r   r^   z)GluonSemantic.histogram.<locals>.<lambda>H  s    7[ r   )r   rF   r9   r   is_intrl   r=   r>   is_boolr7   r   r)   create_histogramr   r3   r   )r    rP   r  r  ra   layout_attrr7   s          r   	histogramzGluonSemantic.histogramB  s    s5;1$&P&PQQQu{!!##%T%TUUUvT!#T#TUUU33D%@@KD%49#++--/[/[\\\;DmmDL11..u|Xt[YY
XJGGGr   worker_num_warpsworker_num_regsc                    t          |          }|t          |          k    sJ d| dt          |           d            |t          |          k    sJ d| dt          |           d            | j        }	|	                                }
|	                                }|	                    |           |                    ||i           }g }|t          |          }|	                    |           d |D             }|	                    |
           t          |          }|		                    |||          
                                                    |                               |           |	                                                    g            |	                    |          }d |D             }t!          |          D ]}t#          ||                   }|	                    |                    |          |          fd	t!          t          |                    D             }t'          |d
 |D                       }|                    ||         |i |           |	                                 |	                                                               fdt!          t          |                    D             }|d S t/          t'          |d |D                                 S )Nzwarp specialize got z partitions but z warp countsz register counts)kwargsc                 6    g | ]}|                                 S r   get_typerW   rs     r   rY   z1GluonSemantic.warp_specialize.<locals>.<listcomp>c  s     ;;;

;;;r   c                 6    g | ]}|                                 S r   r  rW   args     r   rY   z1GluonSemantic.warp_specialize.<locals>.<listcomp>o  s     999S\\^^999r   r   c                 :    g | ]}                     |          S r   )get_argument)rW   jblocks     r   rY   z1GluonSemantic.warp_specialize.<locals>.<listcomp>s  s'    OOOA%,,Q//OOOr   c                     g | ]	}|j         
S r   r\   r  s     r   rY   z1GluonSemantic.warp_specialize.<locals>.<listcomp>t  s    9Z9Z9Zs#(9Z9Z9Zr   )r  caller_contextc                 :    g | ]}                     |          S r   )r   )rW   rL   ws_ops     r   rY   z1GluonSemantic.warp_specialize.<locals>.<listcomp>y  s'    NNN((++NNNr   c                     g | ]	}|j         
S r   r\   r  s     r   rY   z1GluonSemantic.warp_specialize.<locals>.<listcomp>|  s    7X7X7X17X7X7Xr   )rF   r)   get_insertion_point	new_blockset_insertion_point_to_startcall_JitFunctionr   create_warp_yield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_regionr   create_warp_returnset_insertion_point_afterget_operationr   )r    default_argsdefault_partitionworker_argsworker_partitionsr  r  	generatornum_partitionsr)   	insert_ptdefault_blockdefault_resultsmlir_resultsresult_types	mlir_argspartitions_op	arg_typesrL   r!  
block_argsr  r#  s                        @@r   warp_specializezGluonSemantic.warp_specializeN  sR   .//"
 "
 
 
 
e.ee#FVBWBWeee
 
 
 "
 "
 
 
 
h.hh#oBVBVhhh
 
 
 ,//11	  ))++,,];;;#445F]_4``&/@@L!!,///;;l;;; 	''	222(55	..|YHXYY  "",,];;;%%o666 	(()F)F)H)H"MMMAA.QQ99y999	~&& 	) 	)A/:J1:MNNNN44]5M5Ma5P5PR[\\EOOOOs9~~9N9NOOOJ,Z9Z9Zk9Z9Z9Z[[J&&'8';ZPRcq&rrr&&(((())%*=*=*?*?@@@NNNNU3|;L;L5M5MNNN"F(7X7X7X7X7XYYZZZr   )F)/r+   r,   r-   r3   r6   langr   __annotations__r!   r;   r?   r   r.   rO   r   rh   rn   r   rr   rw   r   rl   r   boolr   r   r   r   r   r   r   r   r   r   r   r   r   r   r   staticmethodr   r   r   r  r  rE  __classcell__)rp   s   @r   r0   r0      s        [FD    ' ' '_ _ _49 c     T T T T T T T&5h 58 5 5 5 5 5 5 5Xx XE(H*<$= X X X X X X5X 5U3Z 5H 5 5 5 5 5 5+( +5: +( + + + +  x H      :9 9 9 9 95X 5$s) 5$ 5 5 5 5 5 5
+ + +
1 1 1	+ 	+ 	+ 	+W W W+ + +
G G G
; ; ;+ + +	D 	D 	DD D DY Y Y
 
 
*D D D
& & & X X \X)x'9 ) )"&)+03+?) ) ) )*) 2 )# )UZ[ceh[hUi ) ) ) )$
Hx 
H3 
Hh 
HS[ 
H 
H 
H 
H.[*23-.[JRSV-.[ .[ .[ .[ .[ .[ .[ .[r   r0   )typingr   r   r   r   r   r   triton.language.semanticr    r
   r3   _layoutsr   r   r   triton._C.libtriton.gluon_irr   triton.compiler.code_generatorr   r   r   rG   rH  rJ   r   r   r0   r   r   r   <module>rQ     sf   ; ; ; ; ; ; ; ; ; ; ; ; ; ;  3 3 3 3 3 3       @ @ @ @ @ @ @ @ @ @ 7 7 7 7 7 7 T T T T T T T T7: <F ! ! !xC0 ! ! ! !
	M 	M 	M 	M 	M 	M 	M 	M_[ _[ _[ _[ _[N8, _[ _[ _[ _[ _[r   