o
    h\                     @   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/vscode/kcb/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   box_dim
global_dimglobal_strideelement_stride	elem_typeinterleave_layoutswizzle_mode	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&   int64r   mulr'   r(   r   _determine_swizzle_mode_2d)r   r   r   r   r	   r   element_sizeelement_size_tr   contig_dim_size_in_bytes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)Nr1      @   r   r   r   zblock size too small)r   )r8   r   r   r   r   r5   ]   s   r5   c                 C   s   t | | d S N)r   tensormap_fenceproxy_acquire)r   r   r   r   r   r   h   s   r   r<   )typingr   triton.languager   r   triton._C.libtritonr   __all__dtyper   builtintensorr#   r   	constexprr   r5   r   r   r   r   r   <module>   sN    &