
    h                     d    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y)    )	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	dedee   defd       Z
y	)
r
   baseshapestridesblock_shapelayoutc                 x   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k(  sJ d       t        | j                  t              sJ d       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 aligned   z!Last dimension must be contiguousz Layout must be NVMMASharedLayout)lenr   r   r   r   data_ptrr   r   dtyper   
isinstancer   r	   )selfrank	dtype_str
elem_bytesstrides        e/var/www/html/eduruby.in/venv/lib/python3.12/site-packages/triton/experimental/gluon/nvidia/hopper.py__post_init__zTensorDescriptor.__post_init__   sB   4::4<< D(BOD6*BB(4##$,Fv.FF,ax000xqy666yyy!!#b(A-M/MM-T--.&tyy7	+I6!;
ll3B' 	VFZ'2-2U4UU2	V||B1$I&II$$++'89];]]9    tensorc                 P    t        | | j                  | j                         ||      S )N)r
   r   r   )r#   r   r   s      r    from_tensorzTensorDescriptor.from_tensor    s(    LLMMO
 	
r"   N)__name__
__module____qualname__r   __annotations__r   intr	   r!   staticmethodr%    r"   r    r
   r
   	   s_    
I9#Yc^ 
C 
d3i 
AR 
 
r"   N)dataclassesr   typingr   r   triton._utilsr   r   r   +triton.experimental.gluon.language._layoutsr	   __all__r
   r,   r"   r    <module>r2      s5    !  Z Z I
 
 
 
r"   