o
    cZh\                     @   s   d dl mZ d dlmZ d dlmZ d dlmZ g dZdejfddZ	ej
		dd
ejdejdejdejdejdejfddZej
		dd
ejdejdeej deej dejdejfddZdd Zej
dd
ejdejfddZd	S )    )Sequence)core)semantic)ir)&experimental_device_tensormap_create1d&experimental_device_tensormap_create2d)experimental_tensormap_fenceproxy_acquire
element_tyc                 C   s2   | j dkrdS | j dkrdS | j dkrdS td)N   r                z?element_ty must be a primitive of size 1, 2, or 4 bytes but got)primitive_bitwidth
ValueError)r	    r   [/var/www/auris/lib/python3.10/site-packages/triton/language/extra/cuda/_experimental_tma.py_determine_elem_type   s   


r   Ndesc_ptrglobal_address	load_sizeglobal_size_builderc                 C   sj   t |}t||}t |}t jg dt j|dg}tj| |t||g|gg |t|ddd|d d S )Nr   r   r   r   r   Zbox_dimZ
global_dimglobal_strideelement_strideZ	elem_typeZinterleave_layoutZswizzle_modeZ	fill_modebuilder)r   _constexpr_to_valuer   	to_tensorfullint32tensormap_creater   )r   r   r   r   r	   r   r   r   r   r   r      s"   
	

r   c                    s   t |dksJ t |dksJ dd |D } fdd|D }|jd }tjg |tj d}t||d d }||d  }	|	d	krId	| |d< tjg d
tj d}
tj| | fdd|d d d D |d d d |g|
|
gt	|dt
|	|d d d S )Nr   c                 S   s   g | ]}t |qS r   )r   r   .0xr   r   r   
<listcomp>A   s    z:experimental_device_tensormap_create2d.<locals>.<listcomp>c                       g | ]}t | qS r   r   r   r#   r   r   r   r&   B       r
   r   T   r   c                    r'   r   r(   r#   r   r   r   r&   Q   r)   r   r   )lenr   r   r    Zint64r   mulr!   r"   r   _determine_swizzle_mode_2d)r   r   r   r   r	   r   Zelement_sizeZelement_size_tr   contig_dim_size_in_bytesZelem_strider   r   r   r   6   s0   	

r   c                 C   s,   | dkrdS | dkrdS | dkrdS t d)Nr+      @   r   r   r   zblock size too small)r   )r/   r   r   r   r   r.   ]   s   r.   c                 C   s   t | | d S N)r   Ztensormap_fenceproxy_acquire)r   r   r   r   r   r   h   s   r   r2   )typingr   Ztriton.languager   r   Ztriton._C.libtritonr   __all__Zdtyper   builtinZtensorr   r   Z	constexprr   r.   r   r   r   r   r   <module>   sN    &