Ë
    ŒÓ^jë  ã                   óp   — d dl mZ d dlmZmZ d dlmZmZmZ d dl	m
Z
 d dlmZ dgZe G d„ d«      «       Zy)é    )Ú	dataclass)ÚListÚAny)Úvalidate_block_shapeÚcanonicalize_dtypeÚget_primitive_bitwidth)ÚNVMMASharedLayoutNÚTensorDescriptorc                   óž   — e Zd ZU eed<   ee   ed<   ee   ed<   ee   ed<   eed<   dZe	ed<   dZ
eed	<   d
„ Ze	 	 ddedee   defd„«       Zy)r
   ÚbaseÚshapeÚstridesÚblock_shapeÚlayoutÚzeroÚpaddingFÚround_f32_to_tf32c                 ó²  — t        | j                  «      }t        | j                  «      |k(  s
J d| › «       ‚t        | j                  «      |k(  s
J d| › «       ‚|dkD  sJ d«       ‚|dk  sJ d«       ‚| j                  j                  «       dz  dk(  sJ d«       ‚t        | j                  «       t        | j                  j                  «      }t        |«      dz  }| j                  d d	 D ]  }||z  dz  dk(  rŒJ d
«       ‚ | j                  D ]  }|dkD  rŒ	J d«       ‚ | j                  d	   dk(  sJ d«       ‚t        | j                  t        «      sJ d«       ‚| j                  dk(  s| j                  dk(  sJ d«       ‚| j                  dk(  r'| j                  j                  j                  sJ d«       ‚| j                  r|dk(  sJ d«       ‚|dz  | j                  j                   k(  sJ ‚| j                  j"                  rdnd}| j                  j$                  ||z  z  }| j                  d	   |k\  sJ d|› d| j                  d	   › «       ‚| j                  j"                  rt| j                  j                  «       dz  dk(  sJ d«       ‚| j                  d d	 D ]  }||z  dz  dk(  rŒJ d«       ‚ t&        j(                  j+                  dd«      sJ d«       ‚| j                  j"                  r7| j                  j$                  dk(  sJ d| j                  j$                  › «       ‚| j                  j                   dv sJ d | j                  j                   › «       ‚y )!Nzrank mismatch: r   zrank must not be zeroé   zrank cannot be more than 5é   zbase must be 16-byte alignedé   éÿÿÿÿzstrides must be 16-byte alignedzshape must be positiveé   z!Last dimension must be contiguousz Layout must be NVMMASharedLayoutr   ÚnanzIllegal value for paddingzAPadding option `nan` is only supported for floating point tensorsÚfp32z7round_f32_to_tf32 is only supported for float32 tensorsé   z(Expected block_shape[-1] to be at least z	 but got é    z)For fp4_padded, base must 32-byte alignedz6For fp4_padded, tensor strides must be 32-byte alignedé
   z&fp4_padded requires blackwell or neweré€   zBFP4 padded operands must be swizzled with 128-byte width, but got )r   r   r   z;tensor descriptor dtype must be 8, 16, or 32 bits, but got )Úlenr   r   r   r   Údata_ptrr   r   Údtyper   Ú
isinstancer   r	   r   Úis_floating_pointr   Úelement_bitwidthÚ
fp4_paddedÚswizzle_byte_widthÚtlÚtarget_infoÚcuda_capability_geq)ÚselfÚrankÚ	dtype_strÚ
elem_bytesÚstrideÚ	shape_dimÚpadding_factorÚ	min_blocks           úr/var/www/ramen.bs-engineer-server.com/venv/lib/python3.12/site-packages/triton/experimental/gluon/nvidia/hopper.pyÚ__post_init__zTensorDescriptor.__post_init__   s†  € Ü4—:‘:‹ˆÜ4—<‘<Ó  DÒ(ÐB¨O¸D¸6Ð*BÓBÐ(Ü4×#Ñ#Ó$¨Ò,ÐF°À¸vÐ.FÓFÐ,ØaŠxÐ0Ð0Ó0ˆxØqŠyÐ6Ð6Ó6ˆyØy‰y×!Ñ!Ó# bÑ(¨AÒ-ÐMÐ/MÓMÐ-Ü˜T×-Ñ-Ô.Ü& t§y¡y§¡Ó7ˆ	Ü+¨IÓ6¸!Ñ;ˆ
Ø—l‘l 3 BÐ'ò 	VˆFØ˜ZÑ'¨2Ñ-°Ó2ÐUÐ4UÓUÐ2ð	VàŸ™ò 	;ˆIØ˜q“=Ð:Ð":Ó:=ð	;à|‰|˜BÑ 1Ò$ÐIÐ&IÓIÐ$Ü˜$Ÿ+™+Ô'8Ô9Ð]Ð;]Ó]Ð9Ø|‰|˜vÒ%¨¯©¸Ò)>Ð[Ð@[Ó[Ð>Ø<‰<˜5Ò Ø—9‘9—?‘?×4Ò4ÐyÐ6yÓyÐ4Ø×!Ò!Ø Ò&ÐaÐ(aÓaÐ&Ø˜A‰~ §¡×!=Ñ!=Ò=Ð=Ð=Ø"Ÿk™k×4Ò4™¸!ˆØ—K‘K×2Ñ2°zÀNÑ7RÑSˆ	Ø×Ñ Ñ# yÒ0ð 	bØ6°y°kÀÈ4×K[ÑK[Ð\^ÑK_ÐJ`Ðaó	bÐ0à;‰;×!Ò!Ø—9‘9×%Ñ%Ó'¨"Ñ,°Ò1Ð^Ð3^Ó^Ð1ØŸ,™, s¨Ð+ò qØ Ñ+¨rÑ1°QÓ6ÐpÐ8pÓpÐ6ðqä—>‘>×5Ñ5°b¸!Ô<ÐfÐ>fÓfÐ<Ø—;‘;×)Ò)¨T¯[©[×-KÑ-KÈsÒ-Rð  	zð  WYð  Z^÷  Zeñ  Ze÷  Zxñ  Zxð  Yyð  Uzó  	zÐRØ{‰{×+Ñ+ð 0
ñ 
ð 	hàHÈÏÉ×IeÑIeÐHfÐgó	hñ 
ó    Útensorc           	      óT   — t        | | j                  | j                  «       ||||«      S )N)r
   r   r/   )r6   r   r   r   r   s        r3   Úfrom_tensorzTensorDescriptor.from_tensor8   s0   € ô  ØØL‰LØM‰M‹OØØØØó
ð 	
r5   N)r   F)Ú__name__Ú
__module__Ú__qualname__r   Ú__annotations__r   Úintr	   r   Ústrr   Úboolr4   Ústaticmethodr8   © r5   r3   r
   r
   
   s{   … à
ƒIØ‰9ÓØ#‰YÓØc‘ÓØÓØ€GˆSÓØ#ÐtÓ#ò"hðH Ø\bØ&+ñ

˜Cð 

¨d°3©ið 

ÐARò 

ó ñ

r5   )Údataclassesr   Útypingr   r   Útriton._utilsr   r   r   Ú+triton.experimental.gluon.language._layoutsr	   Útriton.languageÚlanguager(   Ú__all__r
   rA   r5   r3   ú<module>rI      s8   ðÝ !ß ß ZÑ ZÝ IÝ àÐ
€ð ÷8
ð 8
ó ñ8
r5   