Ë
    7^(h\  ã                   ó:  — 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    új/var/www/skyplay_api_hub/venv/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—   € ô ×(Ñ(¨Ó3€IÜ×$Ñ$ [°(Ó;€KÜ×)Ñ)¨*Ó5€JÜ—i‘i  A¤t§z¡z¸HÔEÐF€Nä×ÑØØ%Ü×#Ñ# I¨xÓ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ÖH°q”8×%Ñ% a¨Õ2ÐH€KÐHà×0Ñ0°AÑ5€LÜ—Y‘Y˜r <´·±ÀhÔO€NÜ—L‘L °¸R±À$ÈÓQ€Mà+¨i¸©mÑ;ÐØ #Ò%Ø˜|Ñ+ˆ	�"‰ä—)‘)˜B ¤4§:¡:¸ÔA€Kä×ÑØØ%Ø:CÁDÀbÀD¹/ÖJ°Q”×#Ñ# A xÕ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  ðÝ å  Ý $Ý "ò€ð\ T§Z¡Zó \ð ‡�ð  ñØ�k‰kðà—K‘Kðð �{‰{ðð —‘ð	ð
 —
‘
ðð �j‰jòó ðð8 ‡�ð  ñ#Ø�k‰kð#à—K‘Kð#ð ˜Ÿ™Ñ'ð#ð ˜$Ÿ+™+Ñ&ð	#ð
 —
‘
ð#ð �j‰jò#ó ð#òL1ð ‡�ñ>¸¿¹ð >ÈrÏzÉzò >ó ñ>r   