§
    Ü¶iÕ  ã                   ón   — d dl mZ d dlmZmZ d dlmZ d dlmZm	Z	 dgZ
e G d„ d¦  «        ¦   «         ZdS )é    )Ú	dataclass)ÚListÚAny)Úvalidate_block_shape)ÚPaddedSharedLayoutÚSwizzledSharedLayoutÚTensorDescriptorc                   ó¶   — e Zd ZU eed<   ee         ed<   ee         ed<   ee         ed<   eez  ed<   dZ	e
ed<   d„ Zed	edee         deez  fd
„¦   «         ZdS )r	   ÚbaseÚshapeÚstridesÚblock_shapeÚlayoutÚzeroÚpaddingc                 ó„  — t          | j        ¦  «        }|dk    sJ d|› d¦   «         ‚t          | j        ¦  «        |k    s"J d|› dt          | j        ¦  «        › ¦   «         ‚t          | j        ¦  «        |k    s"J d|› dt          | j        ¦  «        › ¦   «         ‚t	          | j        ¦  «         | j        d         d	k    s
J d
¦   «         ‚t          | j        t          t          f¦  «        s
J d¦   «         ‚t          | j        t          ¦  «        r| j        j	        d	k    s
J d¦   «         ‚| j
        dk    s
J d¦   «         ‚d S )Né   zExpected 2 dimensions but got z dimensionsz	Expected z strides but got zExpected block_shape to have z dimensions but got éÿÿÿÿé   z!Last dimension must be contiguouszBExpected layout to be a PaddedSharedLayout or SwizzledSharedLayoutz3Expected max_phase to be 1 for SwizzledSharedLayoutr   z Only 'zero' padding is supported)Úlenr   r   r   r   Ú
isinstancer   r   r   Ú	max_phaser   )ÚselfÚndims     új/root/projects/openclaw-proxy/venv/lib64/python3.11/site-packages/triton/experimental/gluon/amd/gfx1250.pyÚ__post_init__zTensorDescriptor.__post_init__   so  € Ý4”:‰ŒˆàqŠyˆyˆyÐL¸4ÐLÐLÐL‰yŒyˆyÝ4”<Ñ Ô  DÒ(Ð(Ð(Ð*`°dÐ*`Ð*`ÍSÐQUÔQ]ÑM^ÔM^Ð*`Ð*`Ñ(Ô(Ð(Ý4Ô#Ñ$Ô$¨Ò,Ð,Ð,ØY¨DÐYÐYÅcÈ$Ì,ÑFWÔFWÐYÐYñ -Ô,Ð,å˜TÔ-Ñ.Ô.Ð.ØŒ|˜BÔ 1Ò$Ð$Ð$Ð&IÑ$Ô$Ð$Ý˜$œ+Õ(:Õ<PÐ'QÑRÔRð 	Qð 	QØPñ	Qô 	Qð 	Qåd”kÕ#7Ñ8Ô8ð 	eØ”;Ô(¨AÒ-Ð-Ð-Ð/dÑ-Ô-Ð-ØŒ|˜vÒ%Ð%Ð%Ð'IÑ%Ô%Ð%Ð%Ð%ó    Útensorc                 óV   — t          | | j        |                      ¦   «         ||¦  «        S )a‚   Create a TensorDescriptor object from a tensor.

        Args:
            tensor (torch.Tensor): The input tensor.
            block_shape (List[int]): The block shape of the tensor.
            layout (PaddedSharedLayout | SwizzledSharedLayout): The layout of the tensor in shared memory.

        Returns:
            tensor_descriptor: the created TensorDescriptor object

        )r	   r   Ústride)r   r   r   s      r   Úfrom_tensorzTensorDescriptor.from_tensor!   s%   € õ   ¨¬°f·m²m±o´oÀ{ÐTZÑ[Ô[Ð[r   N)Ú__name__Ú
__module__Ú__qualname__r   Ú__annotations__r   Úintr   r   r   Ústrr   Ústaticmethodr!   © r   r   r	   r	   	   sÌ   € € € € € € à
€I€IIØŒ9ÐÐÑØ#ŒYÐÐÑØc”ÐÐÑØÐ!5Ñ5Ð5Ð5Ñ5Ø€GˆSÐÐÑðJð Jð Jð ð\˜Cð \¨d°3¬ið \ÐASÐVjÑAjð \ð \ð \ñ „\ð\ð \ð \r   N)Údataclassesr   Útypingr   r   Útriton._utilsr   Ú+triton.experimental.gluon.language._layoutsr   r   Ú__all__r	   r)   r   r   ú<module>r/      s©   ðØ !Ð !Ð !Ð !Ð !Ð !Ø Ð Ð Ð Ð Ð Ð Ð Ø .Ð .Ð .Ð .Ð .Ð .Ø `Ð `Ð `Ð `Ð `Ð `Ð `Ð `àÐ
€ð ð$\ð $\ð $\ð $\ð $\ñ $\ô $\ñ „ð$\ð $\ð $\r   