
    nVh\                     :   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Z	ej                  	 dd	ej                  d
ej                  dej                  dej                  dej                  dej                  f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       Zd Zej                  dd	ej                  dej                  fd       Zy)    )Sequence)core)semantic)ir)&experimental_device_tensormap_create1d&experimental_device_tensormap_create2d)experimental_tensormap_fenceproxy_acquire
element_tyc                 x    | j                   dk(  ry| j                   dk(  ry| j                   dk(  ryt        d      )N   r                z?element_ty must be a primitive of size 1, 2, or 4 bytes but got)primitive_bitwidth
ValueError)r
   s    \/home/dcms/DCMS/lib/python3.12/site-packages/triton/language/extra/cuda/_experimental_tma.py_determine_elem_typer      sA    $$)		&	&"	,		&	&"	,Z[[    Ndesc_ptrglobal_address	load_sizeglobal_size_builderc                 R   t        j                  |      }t        j                  ||      }t        j                  |      }t        j                  g dt         j
                  |      g}t        j                  | |t        j                  ||      g|gg |t        |      ddd|       y )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!   s          r   r   r      s     ((3I$$[(;K))*5JiiAtzzHEFN%##Ix89=%&z2r   c                    t        |      dk(  sJ t        |      dk(  sJ |D cg c]  }t        j                  |       }}|D cg c]  }t        j                  ||       }}|j
                  dz  }t        j                  g |t        j                  |      }t        j                  ||d   d|      }	||d   z  }
|
dkD  rd|z  |d<   t        j                  g dt        j                  |      }t        j                  | ||d d d   D cg c]  }t        j                  ||       c}|d d d   |	g||gt        |      dt        |
|      d|	       y c c}w c c}w c c}w )
Nr   r   r   T   r   r   r   )lenr   r'   r   r(   r   r)   int64mulr*   r+   r   _determine_swizzle_mode_2d)r   r   r   r   r
   r   xelement_sizeelement_size_tr    contig_dim_size_in_byteselem_strides               r   r   r   6   s_    y>Q{q   6?@))!,@I@<GHq8%%a2HKH00A5LYYr<hONLLR$QM+im;#%|+	"))B4::AK%:CDbD/JQ##Ax0Jtt$$o#[1&z2/0H)T AH Ks   EEE"c                 <    | dk\  ry| dk\  ry| dk\  ryt        d      )Nr.      @   r   r   r   zblock size too small)r   )r6   r   s     r   r2   r2   ]   s/    3&	!R	'	!R	'/00r   c                 0    t        j                  | |       y N)r   tensormap_fenceproxy_acquire)r   r   s     r   r	   r	   h   s    ))(H=r   r<   )typingr   triton.languager   r   triton._C.libtritonr   __all__dtyper   builtintensorr&   r   	constexprr   r2   r	    r   r   <module>rG      sB      $ "\TZZ \   kkKK {{ 	
 

 jj 8   #kk#KK# '# $++&	#
 

# jj# #L1 > >rzz > >r   