§
    Nb¦i…  ã                   ór   — 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gZe G d„ d¦  «        ¦   «         ZdS )é    )Ú	dataclass)ÚListÚAny)Úvalidate_block_shapeÚcanonicalize_dtypeÚget_primitive_bitwidth)ÚNVMMASharedLayoutÚ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dd	edee         defd
„¦   «         ZdS )r
   ÚbaseÚshapeÚstridesÚblock_shapeÚlayoutÚzeroÚpaddingc                 óf  — 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¦   «         ‚d S 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)Úlenr   r   r   r   Údata_ptrr   r   Údtyper   Ú
isinstancer   r	   r   Úis_floating_point)ÚselfÚrankÚ	dtype_strÚ
elem_bytesÚstrideÚ	shape_dims         úd/root/projects/butler/venv/lib64/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Ð]Ð]Ð;]Ñ]Ô]Ð]ØŒ|˜vÒ%Ð%¨¬¸Ò)>Ð)>Ð)>Ð@[Ñ)>Ô)>Ð)>ØŒ<˜5Ò Ð Ø”9”?Ô4ÐyÐyÐ6yÑyÔyÐyð !Ð ØyÐyó    Útensorc                 óX   — t          | | j        |                      ¦   «         |||¦  «        S )N)r
   r   r#   )r(   r   r   r   s       r%   Úfrom_tensorzTensorDescriptor.from_tensor&   s0   € åØØŒLØMŠM‰OŒOØØØñ
ô 
ð 	
r'   N)r   )Ú__name__Ú
__module__Ú__qualname__r   Ú__annotations__r   Úintr	   r   Ústrr&   Ústaticmethodr*   © r'   r%   r
   r
   	   s¼   € € € € € € à
€I€IIØŒ9ÐÐÑØ#ŒYÐÐÑØc”ÐÐÑØÐÐÑØ€GˆSÐÐÑðzð zð zð( ð
ð 
˜Cð 
¨d°3¬ið 
ÐARð 
ð 
ð 
ñ „\ð
ð 
ð 
r'   N)Údataclassesr   Útypingr   r   Útriton._utilsr   r   r   Ú+triton.experimental.gluon.language._layoutsr	   Ú__all__r
   r2   r'   r%   ú<module>r8      s¥   ðØ !Ð !Ð !Ð !Ð !Ð !Ø Ð Ð Ð Ð Ð Ð Ð Ø ZÐ ZÐ ZÐ ZÐ ZÐ ZÐ ZÐ ZÐ ZÐ ZØ IÐ IÐ IÐ IÐ IÐ IàÐ
€ð ð%
ð %
ð %
ð %
ð %
ñ %
ô %
ñ „ð%
ð %
ð %
r'   