ó
    “~iÕ  ã                   ó`   • S SK Jr  S SKJrJr  S SKJr  S SKJrJ	r	  S/r
\ " S S5      5       rg)é    )Ú	dataclass)ÚListÚAny)Úvalidate_block_shape)ÚPaddedSharedLayoutÚSwizzledSharedLayoutÚTensorDescriptorc                   ó    • \ rS rSr% \\S'   \\   \S'   \\   \S'   \\   \S'   \\	-  \S'   Sr
\\S'   S	 r\S
\S\\   S\\	-  4S j5       rSrg)r	   é	   ÚbaseÚshapeÚstridesÚblock_shapeÚlayoutÚzeroÚpaddingc                 ó˜  • [        U R                  5      nUS:X  d   SU S35       e[        U R                  5      U:X  d    SU S[        U R                  5       35       e[        U R                  5      U:X  d    SU S[        U R                  5       35       e[	        U R                  5        U R                  S   S	:X  d   S
5       e[        U R                  [        [        45      (       d   S5       e[        U R                  [        5      (       a!  U R                  R                  S	:X  d   S5       eU R                  S:X  d   S5       eg )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     Úb/home/mande/repo/quber/.venv/lib/python3.13/site-packages/triton/experimental/gluon/amd/gfx1250.pyÚ__post_init__ÚTensorDescriptor.__post_init__   s;  € Ü�4—:‘:‹ˆà�q‹yÐLÐ:¸4¸&ÀÐLÓLˆyÜ�4—<‘<Ó  DÓ(Ð`¨I°d°VÐ;LÌSÐQU×Q]ÑQ]ÓM^ÐL_Ð*`Ó`Ð(Ü�4×#Ñ#Ó$¨Ó,ð 	ZØ+¨D¨6Ð1EÄcÈ$Ï,É,ÓFWÐEXÐYó	ZÐ,ä˜T×-Ñ-Ô.Ø�|‰|˜BÑ 1Ó$ÐIÐ&IÓIÐ$Ü˜$Ÿ+™+Ô(:Ô<PÐ'Q×RÑRð 	QØPó	QÐRä�d—k‘kÔ#7×8Ñ8Ø—;‘;×(Ñ(¨AÓ-ÐdÐ/dÓdÐ-Ø�|‰|˜vÓ%ÐIÐ'IÓIÑ%ó    Útensorc                 óL   • [        X R                  U R                  5       X5      $ )aI  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_tensorÚTensorDescriptor.from_tensor!   s   € ô   ¯©°f·m±m³oÀ{Ó[Ð[r   © N)Ú__name__Ú
__module__Ú__qualname__Ú__firstlineno__r   Ú__annotations__r   Úintr   r   r   Ústrr   Ústaticmethodr#   Ú__static_attributes__r%   r   r   r	   r	   	   sz   ‡ à
ƒIØ�‰9ÓØ�#‰YÓØ�c‘ÓØÐ!5Ñ5Ó5Ø€GˆSÓò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>r4      s5   ðÝ !ß Ý .ß `àÐ
€ð ÷$\ð $\ó ñ$\r   