o ^i%@sXddlmZddlmZmZddlmZmZmZddl m Z dgZ eGdddZ dS)) dataclass)ListAny)validate_block_shapecanonicalize_dtypeget_primitive_bitwidth)NVMMASharedLayoutTensorDescriptorc@sveZdZUeed<eeed<eeed<eeed<eed<dZe ed<dd Z e dd edeedefd d Z d S)r baseshapestrides block_shapelayoutzeropaddingcCs4t|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 |j dks|j dksJd|j dkr|jjjsJddSdS)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 NVMMASharedLayoutrnanzIllegal value for paddingzAPadding option `nan` is only supported for floating point tensors)lenr r r r data_ptrrrdtyper isinstancerrris_floating_point)selfrank dtype_str elem_bytesstrider!b/var/www/html/RAG/RAG_venv/lib/python3.10/site-packages/triton/experimental/gluon/nvidia/hopper.py __post_init__s$     zTensorDescriptor.__post_init__tensorcCst||j||||S)N)r r r )r$r rrr!r!r" from_tensor$szTensorDescriptor.from_tensorN)r) __name__ __module__ __qualname__r__annotations__rintrrstrr# staticmethodr%r!r!r!r"r s     "N) dataclassesrtypingrr triton._utilsrrr+triton.experimental.gluon.language._layoutsr__all__r r!r!r!r"s