o h@sXddlmZddlmZmZddlmZmZmZddl m Z dgZ eGdddZ dS)) dataclass)ListAny)validate_block_shapecanonicalize_dtypeget_primitive_bitwidth)NVMMASharedLayoutTensorDescriptorc@sheZdZUeed<eeed<eeed<eeed<eed<ddZe dedeedefd d Z d S) r baseshapestrides block_shapelayoutcCst|j}t|j|ksJd|t|j|ks!Jd||dks)Jd|dks1Jd|jddks>Jdt|jt|jj}t |d}|jdd D]}||ddksdJd qV|jd d kspJd t |j t szJd 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_bytesstriderm/home/www/facesmatcher.com/frenv_anti/lib/python3.10/site-packages/triton/experimental/gluon/nvidia/hopper.py __post_init__s    zTensorDescriptor.__post_init__tensorcCst||j|||S)N)r r r)rr rrrr from_tensor szTensorDescriptor.from_tensorN) __name__ __module__ __qualname__r__annotations__rintrr staticmethodrrrrrr s     N) dataclassesrtypingrrZ triton._utilsrrrZ+triton.experimental.gluon.language._layoutsr__all__r rrrrs