a
    khJ6                     @   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eeg ef d	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   Z/var/www/auris/lib/python3.9/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< edddZdd Z	e
e e
e ddd	Zeeed
ddZeeed fddZeeeef d fddZeee ed fddZeee edddZeeed fddZ fddZee
e ed fdd Zd!d" Z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%e&d=d> Z'e(e eeed?f d@dAdBZ)e(e e(e dCdDdEZ*  Z+S )FGluonSemanticbuilderr   c                 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 ]h\}}|| }|dkrT|| q0|dksd||krp|| q0tdt| d t| d t| q0|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 r2 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                      s   d j S Nz=expected expand_dims input to be a distributed_type but got: r   r   r3   r   r   <lambda>6   r:   z+GluonSemantic.expand_dims.<locals>.<lambda>c                      s
   d  S )Nz;expected expand_dims input to have a SliceLayout, but got: r   r   )layoutr   r   r>   9   r:   c                      s   d  dj  S )Nz7expected expand_dims input layout to be sliced in axis z	 but got )dimr   )r4   r?   r   r   r>   ;   r:   )r   insertr)   r   
isinstancer   r   r   r?   r
   r@   r   parentr   Zcreate_expand_dimsr    to_irr!   )r   r3   r4   	dst_shaperet_tyr    r   )r4   r3   r?   r   expand_dims.   s"    



zGluonSemantic.expand_dims)abr5   c                    s8   |  ||\}}t|jg kd t ||}| |S )NzCannot join scalars in gluon)broadcast_impl_valuer   r   superjoinr#   )r   rH   rI   value	__class__r   r   rL   A   s    zGluonSemantic.join)rH   r5   c                    s$   t  |\}}| || |fS r   )rK   splitr#   )r   rH   lhsrhsrN   r   r   rP   G   s    zGluonSemantic.split)r3   dimsr5   c                    s   t  ||}| |S r   )rK   permuter#   )r   r3   rS   rM   rN   r   r   rT   K   s    zGluonSemantic.permute)r3   r   r5   c                    s   t t jtj fdd  j t ttkfdd krR S tD ]F\}}| |krZ|dkrZtd|  d| d| d d	 
qZt jj	 jj
}| j j|| j}| ||S )
Nc                      s   d j S r;   r<   r   r=   r   r   r>   Q   r:   z4GluonSemantic.broadcast_impl_shape.<locals>.<lambda>c                      s   d d  S )Nr&   r'   r   r   )r   	src_shaper   r   r>   S   r:   r   z3Cannot broadcast, the expanded size of the tensor (z ) must match the existing size (z) at non-singleton dimension r(   r'   )r   rB   r   r   r   get_block_shapesr)   r+   r*   r   r?   r   Zcreate_broadcastr    rD   r!   )r   r3   r   r/   itemrF   r    r   )r3   r   rU   r   broadcast_impl_shapeO   s*    

 
z"GluonSemantic.broadcast_impl_shape)rQ   rR   r5   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rtd j
 dj
 | ||}| ||}||fS )Nc                      s
   d S )Nz@expected broadcast left input to be a distributed_type but got: r   r   )lhs_tyr   r   r>   g   r:   z4GluonSemantic.broadcast_impl_value.<locals>.<lambda>c                      s
   d S )NzAexpected broadcast right input to be a distributed_type but got: r   r   )rhs_tyr   r   r>   i   r:   zLayout mismatch in broadcast: z vs )r   Zis_blockrK   rJ   r   rB   r   r   rV   r2   r?   r*   rX   )r   rQ   rR   r$   r%   r.   rN   )rY   rZ   r   rJ   _   s$    

z"GluonSemantic.broadcast_impl_valuec                    s,   || g}t t j||}t j|||dS )N)rF   )r   r   int32rK   arange)r   startendr?   r   rF   rN   r   r   r\   u   s    
zGluonSemantic.arange)r3   rE   can_reorderc                    s&   t | d t |||}| |S )Nz%can_reorder is not supported in gluon)r   rK   reshaper#   )r   r3   rE   r_   rM   rN   r   r   r`   z   s    zGluonSemantic.reshapec                 C   s4   t |j||}| j|| j|j}t ||S r   )r   r   dtyper   Zcreate_splatrD   r    r!   )r   rM   r   r?   rF   r    r   r   r   splat   s    zGluonSemantic.splatc                 C   s   |  ||}| |||S r   )Zmake_scalarrb   )r   r   rM   ra   r?   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                      s
   d S )Nz@expected convert_layout input to be a distributed_type but got: r   r   r"   r   r   r>      r:   z.GluonSemantic.convert_layout.<locals>.<lambda>)r   r   rB   r   r   
element_tyr   r   Zcreate_convert_layoutrD   r    r!   )r   rM   r?   rF   r    r   rd   r   convert_layout   s    
zGluonSemantic.convert_layoutc                 C   sX   t ||||}|d ur2| j|| j|j}n| j|| j}t |||||S r   )r   shared_memory_descriptor_typer   Zcreate_local_allocrD   r    shared_memory_descriptor)r   re   r   r?   rM   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   ra   r   r   Zcreate_local_loadrD   r    r!   )r   mem_descr?   rF   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   rj   rM   r   r   r   shared_store   s    zGluonSemantic.shared_storec                 C   s   | j |j d S r   )r   Zcreate_local_deallocr    )r   rj   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   )r?   r   rg   ra   r   alloc_shaper   Zcreate_memdesc_subviewrD   r    rh   __dict__)r   rj   offsetsr   r?   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   rq   )r   rj   r]   lengthr@   rp   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   rr   rs   rt   r    rq   )r   rj   indexr   rp   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   r7   r/   )rj   r   r   r9      r:   z/GluonSemantic.memdesc_trans.<locals>.<listcomp>c                    s&   g | ]} t  j d  | qS r   )r)   rs   r{   rn   rj   r   r   r9      r:   )re   r   rn   r?   )r)   r   rs   r   rn   r   Zcreate_memdesc_transr    Zget_gluon_layout_from_memdescr   rh   ra   )r   rj   orderr   Znew_alloc_shaper    r?   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   rg   ra   r   rn   r   Zcreate_memdesc_reshaperD   r    rh   ro   )r   rj   r   r?   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   rg   r   Zcreate_memdesc_reinterpretrD   r    rh   ro   )r   rj   ra   r   r?   r"   r    r   r   r   memdesc_reinterpret   s    z!GluonSemantic.memdesc_reinterpretc                 C   s$   |rt |||}n|}| ||S r   )r   r   r!   )r   r8   Z	scalar_tyr.   r?   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                      s   d j S )Nz#expected distributed_type but got: r<   r   )r8   r   r   r>      r:   z2GluonSemantic._check_same_layout.<locals>.<lambda>c                 S   s   g | ]}|j jqS r   )r   r?   r6   r   r   r   r9      r:   z4GluonSemantic._check_same_layout.<locals>.<listcomp>r   c                 3   s   | ]}| kV  qd S r   r   )r7   l)l0r   r   	<genexpr>   r:   z3GluonSemantic._check_same_layout.<locals>.<genexpr>r   c                      s
   d  S )Nz3Expected inputs to have matching layouts, but got: r   r   )layoutsr   r   r>      r:   )r   rB   r   r   r   all)xsr   )r   r   r8   r   _check_same_layout   s    
z GluonSemantic._check_same_layout.)inputsr4   r5   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 sJ d	j	
d
d D  |  sJ 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>      r:   z)GluonSemantic.reduction.<locals>.<lambda>r   c                      s   d d  S )Nz/expected reduction axis to be in the range [0, z
) but got r   r   )r4   rs   r   r   r>      r:   c                    s   g | ]\}}| kr|qS r   r   )r7   r/   s)r4   r   r   r9      r:   z+GluonSemantic.reduction.<locals>.<listcomp>c                 3   s   | ]}|j j kV  qd S r   )r   r   r7   trz   r   r   r      r:   z*GluonSemantic.reduction.<locals>.<genexpr>z-all reduction inputs must have the same shapec                 S   s   g | ]
}|j qS r   )r    r   r   r   r   r9      r:   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
   r?   r   r   Zcreate_reduceverifytuplerange)r   r   r4   Zregion_builder_fnr   )r4   r   rs   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sTJ d| dt | d| j}| }	| }
||
 |j||i d}g }|d urt|}|| dd |D }||	 t|}|	|||
 |
 | | g  ||}dd |D }t|D ]b}||||  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   s   g | ]}|  qS r   get_typer7   rr   r   r   r9     r:   z1GluonSemantic.warp_specialize.<locals>.<listcomp>c                 S   s   g | ]}|  qS r   r   r7   argr   r   r   r9     r:   c                    s   g | ]}  |qS r   )Zget_argument)r7   j)blockr   r   r9     r:   c                 S   s   g | ]
}|j qS r   r<   r   r   r   r   r9     r:   c                    s   g | ]}  |qS r   )r   r{   )ws_opr   r   r9     r:   c                 S   s   g | ]
}|j qS r   r<   r   r   r   r   r9     r:   )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   intr2   r   rG   rL   r   rP   rT   rX   rJ   r\   boolr`   rb   rc   rf   ri   rk   rl   rm   rq   rw   ry   r~   r   r   r   staticmethodr   r   r   r   __classcell__r   r   rN   r   r      sD   

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   