o
    hJ6                     @   s   d dl mZmZmZmZmZ d dlmZ ddlm	Z
 ddl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dZG dd dee ZdS )    )SequenceListTypeVarTupleCallable)TritonSemantic   )_core)SliceLayout)GluonOpBuilder)flatten_values_to_irunflatten_ir_valuesTensorTycondmsg_fnc                 C   s   | s|| d S N )r   r   categoryr   r   r/home/www/facesmatcher.com/frenv_anti/lib/python3.10/site-packages/triton/experimental/gluon/language/_semantic.py_check   s   
r   c                       s  e Zd ZU ejZeZeed< defddZdd Z	de
e de
e fdd	Zd
ededefddZdededef fddZdedeeef f fddZd
edee def fddZd
edee defddZdededef fddZ fdd Zd
ed!e
e d"ef fd#d$Zd%d& Zd'd( Zd)d* Zd+d, Zd-d. Zd/d0 Zd1d2 Zd3d4 Zd5d6 Z d7d8 Z!d9d: Z"d;d< Z#d=d> Z$d?d@ Z%e&dAdB Z'dCe(e dedeedDf fdEdFZ)dGe(e dHe(e fdIdJZ*  Z+S )KGluonSemanticbuilderc                 C   s
   || _ d S r   )r   )selfr   r   r   r   __init__   s   
zGluonSemantic.__init__c                 C   s,   t |jj|j| j|j}| |j|S r   )	ttgldistributed_typetypescalarshaper   Zget_gluon_layout_from_tensorhandletensor)r   r    tyr   r   r   _wrap_tensor_infer_layout   s   z'GluonSemantic._wrap_tensor_infer_layout	lhs_shape	rhs_shapec                 C   s   t |t |krtd| d| g }t|D ]3\}}|| }|dkr*|| q|dks2||kr8|| qtdt| d t| d t| |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#   r$   	ret_shapeileftrightr   r   r   _broadcast_shapes   s*   zGluonSemantic._broadcast_shapesinputaxisreturnc                    s   dd j D }| d  dk r tj 7  ttjtjfdd jjttt	fdd tj
 k fdd tjj|j}| jj || j}| ||S )	Nc                 S   s   g | ]}t |qS r   )r   Z_unwrap_if_constexpr.0xr   r   r   
<listcomp>/       z-GluonSemantic.expand_dims.<locals>.<listcomp>r   r   c                         d j S Nz=expected expand_dims input to be a distributed_type but got: r   r   r2   r   r   <lambda>6       z+GluonSemantic.expand_dims.<locals>.<lambda>c                      
   d  S )Nz;expected expand_dims input to have a SliceLayout, but got: r   r   )layoutr   r   r>   9      
 c                      s   d  dj  S )Nz7expected expand_dims input layout to be sliced in axis z	 but got )dimr   )r3   rA   r   r   r>   ;       )r   insertr(   r   
isinstancer   r   r   rA   r
   rC   r   parentr   Zcreate_expand_dimsr   to_irr    )r   r2   r3   	dst_shaperet_tyr   r   )r3   r2   rA   r   expand_dims.   s"   



zGluonSemantic.expand_dimsabc                    s8   |  ||\}}t|jg kd t ||}| |S )NzCannot join scalars in gluon)broadcast_impl_valuer   r   superjoinr"   )r   rL   rM   value	__class__r   r   rP   A   s   
zGluonSemantic.joinc                    s$   t  |\}}| || |fS r   )rO   splitr"   )r   rL   lhsrhsrR   r   r   rT   G   s   zGluonSemantic.splitdimsc                    s   t  ||}| |S r   )rO   permuter"   )r   r2   rW   rQ   rR   r   r   rX   K   s   
zGluonSemantic.permuter   c                    s   t t jtj fdd  j t ttkfdd kr) S tD ]#\}}| |krP|dkrPtd|  d| d| d d	 
q-t jj	 jj
}| j j|| j}| ||S )
Nc                      r:   r;   r<   r   r=   r   r   r>   Q   r?   z4GluonSemantic.broadcast_impl_shape.<locals>.<lambda>c                         d d  S )Nr%   r&   r   r   )r   	src_shaper   r   r>   S       r   z3Cannot broadcast, the expanded size of the tensor (z ) must match the existing size (z) at non-singleton dimension r'   r&   )r   rF   r   r   r   get_block_shapesr(   r*   r)   r   rA   r   Zcreate_broadcastr   rH   r    )r   r2   r   r.   itemrJ   r   r   )r2   r   rZ   r   broadcast_impl_shapeO   s,   

 z"GluonSemantic.broadcast_impl_shaperU   rV   c                    s   |j  |j   r st ||S tt tj fdd tttjfdd   } }| 	||} j
j
krOtd j
 dj
 | ||}| ||}||fS )Nc                      
   d S )Nz@expected broadcast left input to be a distributed_type but got: r   r   )lhs_tyr   r   r>   g   rB   z4GluonSemantic.broadcast_impl_value.<locals>.<lambda>c                      r_   )NzAexpected broadcast right input to be a distributed_type but got: r   r   )rhs_tyr   r   r>   i   rB   zLayout mismatch in broadcast: z vs )r   Zis_blockrO   rN   r   rF   r   r   r\   r1   rA   r)   r^   )r   rU   rV   r#   r$   r-   rR   )r`   ra   r   rN   _   s$   

z"GluonSemantic.broadcast_impl_valuec                    s,   || g}t t j||}t j|||dS )N)rJ   )r   r   Zint32rO   arange)r   startendrA   r   rJ   rR   r   r   rb   u   s   
zGluonSemantic.arangerI   can_reorderc                    s&   t | d t |||}| |S )Nz%can_reorder is not supported in gluon)r   rO   reshaper"   )r   r2   rI   re   rQ   rR   r   r   rf   z   s   
zGluonSemantic.reshapec                 C   s4   t |j||}| j|| j|j}t ||S r   )r   r   dtyper   Zcreate_splatrH   r   r    )r   rQ   r   rA   rJ   r   r   r   r   splat   s   zGluonSemantic.splatc                 C   s   |  ||}| |||S r   )Zmake_scalarrh   )r   r   rQ   rg   rA   r   r   r   r   full   s   zGluonSemantic.fullc                    sV   |j  tt tj fdd t j j|}| j|	| j|j
}t||S )Nc                      r_   )Nz@expected convert_layout input to be a distributed_type but got: r   r   r!   r   r   r>      rB   z.GluonSemantic.convert_layout.<locals>.<lambda>)r   r   rF   r   r   
element_tyr   r   Zcreate_convert_layoutrH   r   r    )r   rQ   rA   rJ   r   r   rj   r   convert_layout   s   
zGluonSemantic.convert_layoutc                 C   sX   t ||||}|d ur| j|| j|j}n
| j|| j}t |||||S r   )r   shared_memory_descriptor_typer   Zcreate_local_allocrH   r   shared_memory_descriptor)r   rk   r   rA   rQ   r!   r   r   r   r   allocate_shared   s
   zGluonSemantic.allocate_sharedc                 C   s6   t |j|j|}| j|| j|j}t ||S r   )	r   r   rg   r   r   Zcreate_local_loadrH   r   r    )r   mem_descrA   rJ   r   r   r   r   shared_load   s   zGluonSemantic.shared_loadc                 C   s   | j |j|j d S r   )r   Zcreate_local_storer   )r   rp   rQ   r   r   r   shared_store   s   zGluonSemantic.shared_storec                 C   s   | j |j d S r   )r   Zcreate_local_deallocr   )r   rp   r   r   r   shared_dealloc   s   zGluonSemantic.shared_deallocc                 C   sL   |j }t|j|||jj}| j}||||j	|}tj
|fi |jS r   )rA   r   rm   rg   r   alloc_shaper   Zcreate_memdesc_subviewrH   r   rn   __dict__)r   rp   offsetsr   rA   r!   r   r   r   r   r   _memdesc_subview   s
   zGluonSemantic._memdesc_subviewc                 C   sD   | j dg|j }| |j||< t|j}|||< | |||S )Nr   )r   	get_int32rank	to_tensorr   listr   rw   )r   rp   rc   lengthrC   rv   r   r   r   r   memdesc_slice   s
   
zGluonSemantic.memdesc_slicec                 C   s@   |j dd  }| jdg|j }| |j|d< | |||S )Nr   r   )r   r   rx   ry   rz   r   rw   )r   rp   indexr   rv   r   r   r   memdesc_index   s   zGluonSemantic.memdesc_indexc                    s   t |t jksJ dj dt | dfdd|D }jj  d t  j  }| fdd|D 7 }| jj|}| j|}t	j
|j|||dS )Nzsource rank (z) and order length (z) must matchc                    s   g | ]} j | qS r   r   r6   r.   )rp   r   r   r8      r9   z/GluonSemantic.memdesc_trans.<locals>.<listcomp>c                    s&   g | ]} t  j d  | qS r   )r(   ry   r   rt   rp   r   r   r8      s   & )rk   r   rt   rA   )r(   r   ry   r   rt   r   Zcreate_memdesc_transr   Zget_gluon_layout_from_memdescr   rn   rg   )r   rp   orderr   Znew_alloc_shaper   rA   r   r   r   memdesc_trans   s   zGluonSemantic.memdesc_transc                 C   sB   t |j|||jj}| j|| j|j}t j	|fi |j
S r   )r   rm   rg   r   rt   r   Zcreate_memdesc_reshaperH   r   rn   ru   )r   rp   r   rA   r!   r   r   r   r   memdesc_reshape   s   zGluonSemantic.memdesc_reshapec                 C   s<   t ||||}| j|| j|j}t j|fi |jS r   )r   rm   r   Zcreate_memdesc_reinterpretrH   r   rn   ru   )r   rp   rg   r   rA   r!   r   r   r   r   memdesc_reinterpret   s   z!GluonSemantic.memdesc_reinterpretc                 C   s$   |r
t |||}n|}| ||S r   )r   r   r    )r   r7   Z	scalar_tyr-   rA   Zres_tyr   r   r   wrap_tensor   s   zGluonSemantic.wrap_tensorc                    sl   | D ]t tjtjfdd qdd | D d  t t fdddd  D fd	d d S )
Nc                      r:   )Nz#expected distributed_type but got: r<   r   )r7   r   r   r>      r?   z2GluonSemantic._check_same_layout.<locals>.<lambda>c                 S   s   g | ]}|j jqS r   )r   rA   r5   r   r   r   r8          z4GluonSemantic._check_same_layout.<locals>.<listcomp>r   c                 3   s    | ]}| kV  qd S r   r   )r6   l)l0r   r   	<genexpr>   s    z3GluonSemantic._check_same_layout.<locals>.<genexpr>r   c                      r@   )Nz3Expected inputs to have matching layouts, but got: r   r   )layoutsr   r   r>      rB   )r   rF   r   r   r   all)Zxsr   )r   r   r7   r   _check_same_layout   s   
z GluonSemantic._check_same_layoutinputs.c                    s   t  d udd  d jjtt d   kok n   fdd   fddtD t d jjtfddD sOJ d	j	
d
d D  |  seJ tfddttD S )Nc                   S   s   dS )Nz*All-reduce is not yet implemented in gluonr   r   r   r   r   r>      s    z)GluonSemantic.reduction.<locals>.<lambda>r   c                      rY   )Nz/expected reduction axis to be in the range [0, z
) but got r   r   )r3   ry   r   r   r>      r[   c                    s   g | ]
\}}| kr|qS r   r   )r6   r.   s)r3   r   r   r8      s    z+GluonSemantic.reduction.<locals>.<listcomp>c                 3   s    | ]	}|j j kV  qd S r   )r   r   r6   tr   r   r   r      s    z*GluonSemantic.reduction.<locals>.<genexpr>z-all reduction inputs must have the same shapec                 S      g | ]}|j qS r   )r   r   r   r   r   r8      rD   c                 3   s.    | ]} | | jjV  qd S r   )r   
get_resultr   r   r   )r   	reduce_op
ret_layoutr-   r   r   r   r      s
    
)r   r   r   r(   r   r*   r
   rA   r   r   Zcreate_reduceverifytuplerange)r   r   r3   Zregion_builder_fnr   )r3   r   ry   r   r   r-   r   r   r   	reduction   s   (

zGluonSemantic.reductionworker_num_warpsworker_num_regsc                    s  t |}|t |ksJ d| dt | d|t |ks*J d| dt | d| j}| }	| }
||
 |j||i d}g }|d urLt|}|| dd |D }||	 t|}|	|||
 |
 | | g  ||}dd |D }t|D ]0}||||  fd	dtt |D }t|d
d |D }|j|| |i d |  q|  fddtt |D }|d u rd S tt|dd |D S )Nzwarp specialize got z partitions but z warp countsz register counts)kwargsc                 S      g | ]}|  qS r   get_typer6   rr   r   r   r8     r   z1GluonSemantic.warp_specialize.<locals>.<listcomp>c                 S   r   r   r   r6   argr   r   r   r8     r   c                       g | ]}  |qS r   )Zget_argument)r6   j)blockr   r   r8     r9   c                 S   r   r   r<   r   r   r   r   r8     rD   c                    r   r   )r   r   )ws_opr   r   r8     r9   c                 S   r   r   r<   r   r   r   r   r8     rD   )r(   r   Zget_insertion_pointZ	new_blockZset_insertion_point_to_startZcall_JitFunctionr   Zcreate_warp_yieldZrestore_insertion_pointZcreate_warp_specializeZget_default_regionZ	push_backZset_requested_registersZcreate_block_with_parentZget_partition_op_holderZ!create_warp_specialize_partitionsr   Z
get_regionr   Zcreate_warp_returnZset_insertion_point_afterZget_operationr   )r   argsZdefault_partitionZworker_partitionsr   r   	generatorZnum_partitionsr   Z	insert_ptZdefault_blockZdefault_resultsZmlir_resultsZresult_typesZ	mlir_argsZpartitions_opZ	arg_typesr.   Z
block_argsr   )r   r   r   warp_specialize   sP   





zGluonSemantic.warp_specialize),__name__
__module____qualname__r   r    langr   __annotations__r   r"   r   intr1   r   rK   rP   r   rT   rX   r^   rN   rb   boolrf   rh   ri   rl   ro   rq   rr   rs   rw   r}   r   r   r   r   r   staticmethodr   r   r   r   __classcell__r   r   rR   r   r      sF   
 
"
r   N)typingr   r   r   r   r   Ztriton.language.semanticr    r	   r   Z_layoutsr
   Ztriton._C.libtriton.gluon_irr   Ztriton.compiler.code_generatorr   r   r   r)   r   r,   r   r   r   r   r   r   <module>   s    