o i@sTddlmZddlmZmZddlmZddlmZm Z dgZ eGdddZ dS)) dataclass)ListAny)validate_block_shape)PaddedSharedLayoutSwizzledSharedLayoutTensorDescriptorc@s|eZdZUeed<eeed<eeed<eeed<eeBed<dZ e ed<dd Z e d edeedeeBfd d Z d S)rbaseshapestrides block_shapelayoutzeropaddingcCst|j}|dksJd|dt|j|ks%Jd|dt|jt|j|ks9Jd|dt|jt|j|jdd ksIJd t|jttfsUJd t|jtre|jj d kseJd |j d ksnJddS)NzExpected 2 dimensions but got z dimensionsz Expected z strides but got zExpected 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 SwizzledSharedLayoutrz Only 'zero' padding is supported) lenr r r r isinstancer rr max_phaser)selfndimrd/home/ubuntu/Irodori-TTS/.venv/lib/python3.10/site-packages/triton/experimental/gluon/amd/gfx1250.py __post_init__s (  zTensorDescriptor.__post_init__tensorcCst||j|||S)a 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 )rr stride)rr r rrr from_tensor!s zTensorDescriptor.from_tensorN)__name__ __module__ __qualname__r__annotations__rintrrrstrr staticmethodrrrrrr s      $N) dataclassesrtypingrr triton._utilsr+triton.experimental.gluon.language._layoutsrr__all__rrrrrs