§
    OŠtjë  ã                   ó~   — 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dS )é    )Ú	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dS )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k    s
J d¦   «         ‚|dk    s
J d¦   «         ‚| 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    s
J d
¦   «         ‚Œ| j        D ]}|dk    s
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        r|| j                             ¦   «         dz  dk    s
J d¦   «         ‚| j        d d	…         D ]}||z  dz  dk    s
J d¦   «         ‚Œt&          j                             dd¦  «        s
J d¦   «         ‚| j        j        r'| j        j        dk    sJ d| j        j        › �¦   «         ‚| j        j        dv sJ d | j        j        › �¦   «         ‚d S )!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           úe/var/www/html/CA-Chatbot/venv/lib/python3.11/site-packages/triton/experimental/gluon/nvidia/hopper.pyÚ__post_init__zTensorDescriptor.__post_init__   sÐ  € Ý�4”:‰ŒˆÝ�4”<Ñ Ô  DÒ(Ð(Ð(Ð*B¸DÐ*BÐ*BÑ(Ô(Ð(Ý�4Ô#Ñ$Ô$¨Ò,Ð,Ð,Ð.FÀÐ.FÐ.FÑ,Ô,Ð,Ø�aŠxˆxˆxÐ0‰xŒxˆxØ�qŠyˆyˆyÐ6‰yŒyˆyØŒy×!Ò!Ñ#Ô# bÑ(¨AÒ-Ð-Ð-Ð/MÑ-Ô-Ð-Ý˜TÔ-Ñ.Ô.Ð.Ý& t¤y¤Ñ7Ô7ˆ	Ý+¨IÑ6Ô6¸!Ñ;ˆ
Ø”l 3 B 3Ô'ð 	Vð 	VˆFØ˜ZÑ'¨2Ñ-°Ò2Ð2Ð2Ð4UÑ2Ô2Ð2Ð2Øœð 	;ð 	;ˆIØ˜q’=�=�=Ð":‘=”=�=�=ØŒ|˜BÔ 1Ò$Ð$Ð$Ð&IÑ$Ô$Ð$Ý˜$œ+Õ'8Ñ9Ô9Ð]Ð]Ð;]Ñ]Ô]Ð9ØŒ|˜vÒ%Ð%¨¬¸Ò)>Ð)>Ð)>Ð@[Ñ)>Ô)>Ð>ØŒ<˜5Ò Ð Ø”9”?Ô4ÐyÐyÐ6yÑyÔyÐ4ØÔ!ð 	bØ Ò&Ð&Ð&Ð(aÑ&Ô&Ð&Ø˜A‰~ ¤Ô!=Ò=Ð=Ð=Ð=Ø"œkÔ4Ð;˜˜¸!ˆØ”KÔ2°zÀNÑ7RÑSˆ	ØÔ Ô# yÒ0Ð0Ð0Øa°yÐaÐaÈ4ÔK[Ð\^ÔK_ÐaÐañ 1Ô0Ð0àŒ;Ô!ð 	gØ”9×%Ò%Ñ'Ô'¨"Ñ,°Ò1Ð1Ð1Ð3^Ñ1Ô1Ð1Øœ, s¨ sÔ+ð qð q�Ø Ñ+¨rÑ1°QÒ6Ð6Ð6Ð8pÑ6Ô6Ð6Ð6Ý”>×5Ò5°b¸!Ñ<Ô<ÐfÐfÐ>fÑfÔfÐ<Ø”;Ô)ð  	z¨T¬[Ô-KÈsÒ-RÐ-RÐ-Rð  Uzð  Z^ô  Zeô  Zxð  Uzð  UzÑ-RÔ-RÐRØŒ{Ô+ð 0
ð 
ð 
ð 
àgÈÌÔIeÐgÐgñ
ô 
ð 
ð 
ð 
ó    Útensorc           	      óZ   — t          | | j        |                      ¦   «         ||||¦  «        S )N)r
   r   r/   )r6   r   r   r   r   s        r3   Úfrom_tensorzTensorDescriptor.from_tensor8   s5   € õ  ØØŒLØ�MŠM‰OŒOØØØØñ
ô 
ð 	
r5   N)r   F)Ú__name__Ú
__module__Ú__qualname__r   Ú__annotations__r   Úintr	   r   Ústrr   Úboolr4   Ústaticmethodr8   © r5   r3   r
   r
   
   sÔ   € € € € € € à
€I€I�IØ�Œ9ÐÐÑØ�#ŒYÐÐÑØ�c”ÐÐÑØÐÐÑØ€GˆSÐÐÑØ#Ð�tÐ#Ð#Ñ#ð"hð "hð "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      s·   ðØ !Ð !Ð !Ð !Ð !Ð !Ø Ð Ð Ð Ð Ð Ð Ð Ø ZÐ ZÐ ZÐ ZÐ ZÐ ZÐ ZÐ ZÐ ZÐ ZØ IÐ IÐ IÐ IÐ IÐ IØ Ð Ð Ð Ð Ð àÐ
€ð ð8
ð 8
ð 8
ð 8
ð 8
ñ 8
ô 8
ñ „ð8
ð 8
ð 8
r5   