a kh@sXddlmZddlmZmZddlmZmZmZddl m Z dgZ eGdddZ dS)) dataclass)ListAny)validate_block_shapecanonicalize_dtypeget_primitive_bitwidth)NVMMASharedLayoutTensorDescriptorc@sdeZdZUeed<eeed<eeed<eeed<eed<ddZe eeeedd d Z d S) r baseshapestrides block_shapelayoutcCst|j}t|j|ks&Jd|t|j|ksBJd||dksRJd|dksbJd|jddks|Jdt|jt|jj}t |d}|jdd D]}||ddksJd q|jd d ksJd t |j t sJd dS)Nzrank mismatch: rzrank must not be zerozrank cannot be more than 5zbase must be 16-byte alignedzstrides must be 16-byte alignedz!Last dimension must be contiguousz Layout must be NVMMASharedLayout) lenr r r r Zdata_ptrrrZdtyper isinstancerr)selfZrankZ dtype_strZ elem_bytesstriderU/var/www/auris/lib/python3.9/site-packages/triton/experimental/gluon/nvidia/hopper.py __post_init__s    zTensorDescriptor.__post_init__Ztensorr rcCst||j|||S)N)r r rrrrr from_tensor szTensorDescriptor.from_tensorN) __name__ __module__ __qualname__r__annotations__rintrr staticmethodrrrrrr s    N) Z dataclassesrtypingrrZ triton._utilsrrrZ+triton.experimental.gluon.language._layoutsr__all__r rrrrs