# Autogenerated by mlir-tblgen; don't manually edit.

from enum import IntEnum, auto, IntFlag
from ._ods_common import _cext as _ods_cext
from ..ir import register_attribute_builder
_ods_ir = _ods_cext.ir

class BarrierReduction(IntEnum):
    """NVVM barrier reduction operation"""

    POPC = 0
    AND = 1
    OR = 2

    def __str__(self):
        if self is BarrierReduction.POPC:
            return "popc"
        if self is BarrierReduction.AND:
            return "and"
        if self is BarrierReduction.OR:
            return "or"
        raise ValueError("Unknown BarrierReduction enum entry.")



@register_attribute_builder("BarrierReduction")
def _barrierreduction(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

class BlockScaleFormat(IntEnum):
    """MMA Block Scale Format"""

    UE8M0 = 0
    UE4M3 = 1

    def __str__(self):
        if self is BlockScaleFormat.UE8M0:
            return "ue8m0"
        if self is BlockScaleFormat.UE4M3:
            return "ue4m3"
        raise ValueError("Unknown BlockScaleFormat enum entry.")



@register_attribute_builder("BlockScaleFormat")
def _blockscaleformat(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

class CTAGroupKind(IntEnum):
    """NVVM CTA group kind"""

    CTA_1 = 0
    CTA_2 = 1

    def __str__(self):
        if self is CTAGroupKind.CTA_1:
            return "cta_1"
        if self is CTAGroupKind.CTA_2:
            return "cta_2"
        raise ValueError("Unknown CTAGroupKind enum entry.")



@register_attribute_builder("CTAGroupKind")
def _ctagroupkind(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

class CacheEvictionPriority(IntEnum):
    """NVVM Cache Eviction Priority"""

    EvictNormal = 0
    EvictFirst = 1
    EvictLast = 2
    EvictUnchanged = 3
    NoAllocate = 4

    def __str__(self):
        if self is CacheEvictionPriority.EvictNormal:
            return "evict_normal"
        if self is CacheEvictionPriority.EvictFirst:
            return "evict_first"
        if self is CacheEvictionPriority.EvictLast:
            return "evict_last"
        if self is CacheEvictionPriority.EvictUnchanged:
            return "evict_unchanged"
        if self is CacheEvictionPriority.NoAllocate:
            return "no_allocate"
        raise ValueError("Unknown CacheEvictionPriority enum entry.")



class ClusterLaunchControlQueryType(IntEnum):
    """NVVM ClusterLaunchControlQueryType"""

    IS_CANCELED = 0
    GET_FIRST_CTA_ID_X = 1
    GET_FIRST_CTA_ID_Y = 2
    GET_FIRST_CTA_ID_Z = 3

    def __str__(self):
        if self is ClusterLaunchControlQueryType.IS_CANCELED:
            return "is_canceled"
        if self is ClusterLaunchControlQueryType.GET_FIRST_CTA_ID_X:
            return "get_first_cta_id_x"
        if self is ClusterLaunchControlQueryType.GET_FIRST_CTA_ID_Y:
            return "get_first_cta_id_y"
        if self is ClusterLaunchControlQueryType.GET_FIRST_CTA_ID_Z:
            return "get_first_cta_id_z"
        raise ValueError("Unknown ClusterLaunchControlQueryType enum entry.")



class DotAccumulateType(IntEnum):
    """NVVM DotAccumulateType"""

    SIGNED = 1
    UNSIGNED = 0

    def __str__(self):
        if self is DotAccumulateType.SIGNED:
            return "signed"
        if self is DotAccumulateType.UNSIGNED:
            return "unsigned"
        raise ValueError("Unknown DotAccumulateType enum entry.")



@register_attribute_builder("DotAccumulateType")
def _dotaccumulatetype(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

class FPRoundingMode(IntEnum):
    """NVVM FPRoundingMode kind"""

    NONE = 0
    RN = 1
    RM = 2
    RP = 3
    RZ = 4
    RNA = 5
    RS = 6

    def __str__(self):
        if self is FPRoundingMode.NONE:
            return "none"
        if self is FPRoundingMode.RN:
            return "rn"
        if self is FPRoundingMode.RM:
            return "rm"
        if self is FPRoundingMode.RP:
            return "rp"
        if self is FPRoundingMode.RZ:
            return "rz"
        if self is FPRoundingMode.RNA:
            return "rna"
        if self is FPRoundingMode.RS:
            return "rs"
        raise ValueError("Unknown FPRoundingMode enum entry.")



@register_attribute_builder("FPRoundingMode")
def _fproundingmode(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

class GridDepActionKind(IntEnum):
    """Action kind for grid dependency control"""

    wait = 0
    launch_dependents = 1

    def __str__(self):
        if self is GridDepActionKind.wait:
            return "wait"
        if self is GridDepActionKind.launch_dependents:
            return "launch_dependents"
        raise ValueError("Unknown GridDepActionKind enum entry.")



class LdStMatrixEltType(IntEnum):
    """Element type for ldmatrix and stmatrix"""

    B16 = 0
    B8 = 1
    B8X16_B6X16_P32 = 2
    B8X16_B4X16_P64 = 3

    def __str__(self):
        if self is LdStMatrixEltType.B16:
            return "b16"
        if self is LdStMatrixEltType.B8:
            return "b8"
        if self is LdStMatrixEltType.B8X16_B6X16_P32:
            return "b8x16.b6x16_p32"
        if self is LdStMatrixEltType.B8X16_B4X16_P64:
            return "b8x16.b4x16_p64"
        raise ValueError("Unknown LdStMatrixEltType enum entry.")



@register_attribute_builder("LdStMatrixEltType")
def _ldstmatrixelttype(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

class LoadCacheModifierKind(IntEnum):
    """NVVM load cache modifier kind"""

    CA = 0
    CG = 1
    CS = 2
    LU = 3
    CV = 4

    def __str__(self):
        if self is LoadCacheModifierKind.CA:
            return "ca"
        if self is LoadCacheModifierKind.CG:
            return "cg"
        if self is LoadCacheModifierKind.CS:
            return "cs"
        if self is LoadCacheModifierKind.LU:
            return "lu"
        if self is LoadCacheModifierKind.CV:
            return "cv"
        raise ValueError("Unknown LoadCacheModifierKind enum entry.")



@register_attribute_builder("LoadCacheModifierKind")
def _loadcachemodifierkind(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

class MMAB1Op(IntEnum):
    """MMA binary operations"""

    none = 0
    xor_popc = 1
    and_popc = 2

    def __str__(self):
        if self is MMAB1Op.none:
            return "none"
        if self is MMAB1Op.xor_popc:
            return "xor_popc"
        if self is MMAB1Op.and_popc:
            return "and_popc"
        raise ValueError("Unknown MMAB1Op enum entry.")



@register_attribute_builder("MMAB1Op")
def _mmab1op(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

class MMABlockScaleKind(IntEnum):
    """Block Scale Kind"""

    MXF8F6F4 = 0
    MXF4 = 1
    MXF4NVF4 = 2

    def __str__(self):
        if self is MMABlockScaleKind.MXF8F6F4:
            return "mxf8f6f4"
        if self is MMABlockScaleKind.MXF4:
            return "mxf4"
        if self is MMABlockScaleKind.MXF4NVF4:
            return "mxf4nvf4"
        raise ValueError("Unknown MMABlockScaleKind enum entry.")



@register_attribute_builder("MMABlockScaleKind")
def _mmablockscalekind(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

class MMAFrag(IntEnum):
    """NVVM MMA frag type"""

    a = 0
    b = 1
    c = 2

    def __str__(self):
        if self is MMAFrag.a:
            return "a"
        if self is MMAFrag.b:
            return "b"
        if self is MMAFrag.c:
            return "c"
        raise ValueError("Unknown MMAFrag enum entry.")



@register_attribute_builder("MMAFrag")
def _mmafrag(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

class MMAIntOverflow(IntEnum):
    """MMA overflow options"""

    satfinite = 1
    wrapped = 0

    def __str__(self):
        if self is MMAIntOverflow.satfinite:
            return "satfinite"
        if self is MMAIntOverflow.wrapped:
            return "wrapped"
        raise ValueError("Unknown MMAIntOverflow enum entry.")



@register_attribute_builder("MMAIntOverflow")
def _mmaintoverflow(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

class MMAKind(IntEnum):
    """MMA operation kind"""

    f8f6f4 = 0

    def __str__(self):
        if self is MMAKind.f8f6f4:
            return "f8f6f4"
        raise ValueError("Unknown MMAKind enum entry.")



@register_attribute_builder("MMAKind")
def _mmakind(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

class MMALayout(IntEnum):
    """NVVM MMA layout"""

    row = 0
    col = 1

    def __str__(self):
        if self is MMALayout.row:
            return "row"
        if self is MMALayout.col:
            return "col"
        raise ValueError("Unknown MMALayout enum entry.")



@register_attribute_builder("MMALayout")
def _mmalayout(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

class MMATypes(IntEnum):
    """NVVM MMA types"""

    f16 = 0
    f32 = 1
    tf32 = 2
    bf16 = 9
    s8 = 4
    u8 = 3
    s32 = 5
    s4 = 8
    u4 = 7
    b1 = 6
    f64 = 10
    e4m3 = 11
    e5m2 = 12
    e3m2 = 13
    e2m3 = 14
    e2m1 = 15

    def __str__(self):
        if self is MMATypes.f16:
            return "f16"
        if self is MMATypes.f32:
            return "f32"
        if self is MMATypes.tf32:
            return "tf32"
        if self is MMATypes.bf16:
            return "bf16"
        if self is MMATypes.s8:
            return "s8"
        if self is MMATypes.u8:
            return "u8"
        if self is MMATypes.s32:
            return "s32"
        if self is MMATypes.s4:
            return "s4"
        if self is MMATypes.u4:
            return "u4"
        if self is MMATypes.b1:
            return "b1"
        if self is MMATypes.f64:
            return "f64"
        if self is MMATypes.e4m3:
            return "e4m3"
        if self is MMATypes.e5m2:
            return "e5m2"
        if self is MMATypes.e3m2:
            return "e3m2"
        if self is MMATypes.e2m3:
            return "e2m3"
        if self is MMATypes.e2m1:
            return "e2m1"
        raise ValueError("Unknown MMATypes enum entry.")



@register_attribute_builder("MMATypes")
def _mmatypes(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

class MatchSyncKind(IntEnum):
    """NVVM match sync kind"""

    any = 0
    all = 1

    def __str__(self):
        if self is MatchSyncKind.any:
            return "any"
        if self is MatchSyncKind.all:
            return "all"
        raise ValueError("Unknown MatchSyncKind enum entry.")



@register_attribute_builder("MatchSyncKind")
def _matchsynckind(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

class MemOrderKind(IntEnum):
    """NVVM Memory Ordering kind"""

    WEAK = 0
    RELAXED = 1
    ACQUIRE = 2
    RELEASE = 3
    ACQ_REL = 4
    SC = 5
    MMIO = 6
    VOLATILE = 7

    def __str__(self):
        if self is MemOrderKind.WEAK:
            return "weak"
        if self is MemOrderKind.RELAXED:
            return "relaxed"
        if self is MemOrderKind.ACQUIRE:
            return "acquire"
        if self is MemOrderKind.RELEASE:
            return "release"
        if self is MemOrderKind.ACQ_REL:
            return "acq_rel"
        if self is MemOrderKind.SC:
            return "sc"
        if self is MemOrderKind.MMIO:
            return "mmio"
        if self is MemOrderKind.VOLATILE:
            return "volatile"
        raise ValueError("Unknown MemOrderKind enum entry.")



@register_attribute_builder("MemOrderKind")
def _memorderkind(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

class MemScopeKind(IntEnum):
    """NVVM Memory Scope kind"""

    CTA = 0
    CLUSTER = 1
    GPU = 2
    SYS = 3

    def __str__(self):
        if self is MemScopeKind.CTA:
            return "cta"
        if self is MemScopeKind.CLUSTER:
            return "cluster"
        if self is MemScopeKind.GPU:
            return "gpu"
        if self is MemScopeKind.SYS:
            return "sys"
        raise ValueError("Unknown MemScopeKind enum entry.")



@register_attribute_builder("MemScopeKind")
def _memscopekind(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

class NVVMMemorySpace(IntEnum):
    """NVVM Memory Space"""

    Generic = 0
    Global = 1
    Shared = 3
    Constant = 4
    Local = 5
    Tensor = 6
    SharedCluster = 7

    def __str__(self):
        if self is NVVMMemorySpace.Generic:
            return "generic"
        if self is NVVMMemorySpace.Global:
            return "global"
        if self is NVVMMemorySpace.Shared:
            return "shared"
        if self is NVVMMemorySpace.Constant:
            return "constant"
        if self is NVVMMemorySpace.Local:
            return "local"
        if self is NVVMMemorySpace.Tensor:
            return "tensor"
        if self is NVVMMemorySpace.SharedCluster:
            return "shared_cluster"
        raise ValueError("Unknown NVVMMemorySpace enum entry.")



class PermuteMode(IntEnum):
    """NVVM permute mode"""

    DEFAULT = 0
    F4E = 1
    B4E = 2
    RC8 = 3
    ECL = 4
    ECR = 5
    RC16 = 6

    def __str__(self):
        if self is PermuteMode.DEFAULT:
            return "default"
        if self is PermuteMode.F4E:
            return "f4e"
        if self is PermuteMode.B4E:
            return "b4e"
        if self is PermuteMode.RC8:
            return "rc8"
        if self is PermuteMode.ECL:
            return "ecl"
        if self is PermuteMode.ECR:
            return "ecr"
        if self is PermuteMode.RC16:
            return "rc16"
        raise ValueError("Unknown PermuteMode enum entry.")



@register_attribute_builder("PermuteMode")
def _permutemode(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

class PrefetchCacheLevel(IntEnum):
    """NVVM Prefetch Cache Level"""

    L1 = 0
    L2 = 1

    def __str__(self):
        if self is PrefetchCacheLevel.L1:
            return "L1"
        if self is PrefetchCacheLevel.L2:
            return "L2"
        raise ValueError("Unknown PrefetchCacheLevel enum entry.")



class ProxyKind(IntEnum):
    """Proxy kind"""

    alias = 0
    async_ = 1
    async_global = 2
    async_shared = 3
    TENSORMAP = 4
    GENERIC = 5

    def __str__(self):
        if self is ProxyKind.alias:
            return "alias"
        if self is ProxyKind.async_:
            return "async"
        if self is ProxyKind.async_global:
            return "async.global"
        if self is ProxyKind.async_shared:
            return "async.shared"
        if self is ProxyKind.TENSORMAP:
            return "tensormap"
        if self is ProxyKind.GENERIC:
            return "generic"
        raise ValueError("Unknown ProxyKind enum entry.")



@register_attribute_builder("ProxyKind")
def _proxykind(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

class ReduxKind(IntEnum):
    """NVVM redux kind"""

    ADD = 1
    AND = 2
    MAX = 3
    MIN = 4
    OR = 5
    UMAX = 6
    UMIN = 7
    XOR = 8
    FMIN = 9
    FMAX = 10

    def __str__(self):
        if self is ReduxKind.ADD:
            return "add"
        if self is ReduxKind.AND:
            return "and"
        if self is ReduxKind.MAX:
            return "max"
        if self is ReduxKind.MIN:
            return "min"
        if self is ReduxKind.OR:
            return "or"
        if self is ReduxKind.UMAX:
            return "umax"
        if self is ReduxKind.UMIN:
            return "umin"
        if self is ReduxKind.XOR:
            return "xor"
        if self is ReduxKind.FMIN:
            return "fmin"
        if self is ReduxKind.FMAX:
            return "fmax"
        raise ValueError("Unknown ReduxKind enum entry.")



@register_attribute_builder("ReduxKind")
def _reduxkind(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

class SaturationMode(IntEnum):
    """NVVM SaturationMode kind"""

    NONE = 0
    SATFINITE = 1

    def __str__(self):
        if self is SaturationMode.NONE:
            return "none"
        if self is SaturationMode.SATFINITE:
            return "satfinite"
        raise ValueError("Unknown SaturationMode enum entry.")



@register_attribute_builder("SaturationMode")
def _saturationmode(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

class ScaleVecSize(IntEnum):
    """MMA Scale Vector Sizes"""

    X1 = 0
    X2 = 1
    X4 = 2

    def __str__(self):
        if self is ScaleVecSize.X1:
            return "x1"
        if self is ScaleVecSize.X2:
            return "x2"
        if self is ScaleVecSize.X4:
            return "x4"
        raise ValueError("Unknown ScaleVecSize enum entry.")



@register_attribute_builder("ScaleVecSize")
def _scalevecsize(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

class SetMaxRegisterAction(IntEnum):
    """NVVM set max register action"""

    decrease = 1
    increase = 0

    def __str__(self):
        if self is SetMaxRegisterAction.decrease:
            return "decrease"
        if self is SetMaxRegisterAction.increase:
            return "increase"
        raise ValueError("Unknown SetMaxRegisterAction enum entry.")



@register_attribute_builder("SetMaxRegisterAction")
def _setmaxregisteraction(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

class SharedSpace(IntEnum):
    """Shared memory space"""

    shared_cta = 0
    shared_cluster = 1

    def __str__(self):
        if self is SharedSpace.shared_cta:
            return "cta"
        if self is SharedSpace.shared_cluster:
            return "cluster"
        raise ValueError("Unknown SharedSpace enum entry.")



@register_attribute_builder("SharedSpace")
def _sharedspace(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

class ShflKind(IntEnum):
    """NVVM shuffle kind"""

    bfly = 0
    up = 1
    down = 2
    idx = 3

    def __str__(self):
        if self is ShflKind.bfly:
            return "bfly"
        if self is ShflKind.up:
            return "up"
        if self is ShflKind.down:
            return "down"
        if self is ShflKind.idx:
            return "idx"
        raise ValueError("Unknown ShflKind enum entry.")



@register_attribute_builder("ShflKind")
def _shflkind(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

class TMALoadMode(IntEnum):
    """NVVM TMA Load Mode"""

    TILE = 0
    IM2COL = 1
    IM2COL_W = 2
    IM2COL_W_128 = 3
    TILE_GATHER4 = 4

    def __str__(self):
        if self is TMALoadMode.TILE:
            return "tile"
        if self is TMALoadMode.IM2COL:
            return "im2col"
        if self is TMALoadMode.IM2COL_W:
            return "im2col_w"
        if self is TMALoadMode.IM2COL_W_128:
            return "im2col_w_128"
        if self is TMALoadMode.TILE_GATHER4:
            return "tile_gather4"
        raise ValueError("Unknown TMALoadMode enum entry.")



@register_attribute_builder("TMALoadMode")
def _tmaloadmode(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

class TMAReduxKind(IntEnum):
    """NVVM TMA redux kind"""

    ADD = 0
    MAX = 2
    MIN = 1
    INC = 3
    DEC = 4
    AND = 5
    OR = 6
    XOR = 7

    def __str__(self):
        if self is TMAReduxKind.ADD:
            return "add"
        if self is TMAReduxKind.MAX:
            return "max"
        if self is TMAReduxKind.MIN:
            return "min"
        if self is TMAReduxKind.INC:
            return "inc"
        if self is TMAReduxKind.DEC:
            return "dec"
        if self is TMAReduxKind.AND:
            return "and"
        if self is TMAReduxKind.OR:
            return "or"
        if self is TMAReduxKind.XOR:
            return "xor"
        raise ValueError("Unknown TMAReduxKind enum entry.")



@register_attribute_builder("TMAReduxKind")
def _tmareduxkind(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

class TMAStoreMode(IntEnum):
    """NVVM TMA Store Mode"""

    TILE = 0
    IM2COL = 1
    TILE_SCATTER4 = 2

    def __str__(self):
        if self is TMAStoreMode.TILE:
            return "tile"
        if self is TMAStoreMode.IM2COL:
            return "im2col"
        if self is TMAStoreMode.TILE_SCATTER4:
            return "tile_scatter4"
        raise ValueError("Unknown TMAStoreMode enum entry.")



@register_attribute_builder("TMAStoreMode")
def _tmastoremode(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

class Tcgen05CpMulticast(IntEnum):
    """tcgen05 cp multicast"""

    NONE = 0
    WARPX2_02_13 = 1
    WARPX2_01_23 = 2
    WARPX4 = 3

    def __str__(self):
        if self is Tcgen05CpMulticast.NONE:
            return "none"
        if self is Tcgen05CpMulticast.WARPX2_02_13:
            return "warpx2_02_13"
        if self is Tcgen05CpMulticast.WARPX2_01_23:
            return "warpx2_01_23"
        if self is Tcgen05CpMulticast.WARPX4:
            return "warpx4"
        raise ValueError("Unknown Tcgen05CpMulticast enum entry.")



@register_attribute_builder("Tcgen05CpMulticast")
def _tcgen05cpmulticast(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

class Tcgen05CpShape(IntEnum):
    """tcgen05 cp shapes"""

    SHAPE_128x256b = 0
    SHAPE_4x256b = 1
    SHAPE_128x128b = 2
    SHAPE_64x128b = 3
    SHAPE_32x128b = 4

    def __str__(self):
        if self is Tcgen05CpShape.SHAPE_128x256b:
            return "shape_128x256b"
        if self is Tcgen05CpShape.SHAPE_4x256b:
            return "shape_4x256b"
        if self is Tcgen05CpShape.SHAPE_128x128b:
            return "shape_128x128b"
        if self is Tcgen05CpShape.SHAPE_64x128b:
            return "shape_64x128b"
        if self is Tcgen05CpShape.SHAPE_32x128b:
            return "shape_32x128b"
        raise ValueError("Unknown Tcgen05CpShape enum entry.")



@register_attribute_builder("Tcgen05CpShape")
def _tcgen05cpshape(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

class Tcgen05CpSrcFormat(IntEnum):
    """tcgen05 cp source format"""

    B6x16_P32 = 0
    B4x16_P64 = 1

    def __str__(self):
        if self is Tcgen05CpSrcFormat.B6x16_P32:
            return "b6x16_p32"
        if self is Tcgen05CpSrcFormat.B4x16_P64:
            return "b4x16_p64"
        raise ValueError("Unknown Tcgen05CpSrcFormat enum entry.")



@register_attribute_builder("Tcgen05CpSrcFormat")
def _tcgen05cpsrcformat(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

class Tcgen05FenceKind(IntEnum):
    """NVVM Tcgen05 fence kind"""

    BEFORE_THREAD_SYNC = 0
    AFTER_THREAD_SYNC = 1

    def __str__(self):
        if self is Tcgen05FenceKind.BEFORE_THREAD_SYNC:
            return "before"
        if self is Tcgen05FenceKind.AFTER_THREAD_SYNC:
            return "after"
        raise ValueError("Unknown Tcgen05FenceKind enum entry.")



@register_attribute_builder("Tcgen05FenceKind")
def _tcgen05fencekind(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

class Tcgen05LdStShape(IntEnum):
    """allowed 32-bit signless integer cases: 0, 1, 2, 3, 4"""

    SHAPE_16X64B = 0
    SHAPE_16X128B = 1
    SHAPE_16X256B = 2
    SHAPE_32X32B = 3
    SHAPE_16X32BX2 = 4

    def __str__(self):
        if self is Tcgen05LdStShape.SHAPE_16X64B:
            return "shape_16x64b"
        if self is Tcgen05LdStShape.SHAPE_16X128B:
            return "shape_16x128b"
        if self is Tcgen05LdStShape.SHAPE_16X256B:
            return "shape_16x256b"
        if self is Tcgen05LdStShape.SHAPE_32X32B:
            return "shape_32x32b"
        if self is Tcgen05LdStShape.SHAPE_16X32BX2:
            return "shape_16x32bx2"
        raise ValueError("Unknown Tcgen05LdStShape enum entry.")



@register_attribute_builder("Tcgen05LdStShape")
def _tcgen05ldstshape(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

class Tcgen05MMABlockScale(IntEnum):
    """tcgen05.mma block scale attribute"""

    DEFAULT = 0
    BLOCK16 = 1
    BLOCK32 = 2

    def __str__(self):
        if self is Tcgen05MMABlockScale.DEFAULT:
            return "default"
        if self is Tcgen05MMABlockScale.BLOCK16:
            return "block16"
        if self is Tcgen05MMABlockScale.BLOCK32:
            return "block32"
        raise ValueError("Unknown Tcgen05MMABlockScale enum entry.")



@register_attribute_builder("Tcgen05MMABlockScale")
def _tcgen05mmablockscale(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

class Tcgen05MMACollectorBBuffer(IntEnum):
    """tcgen05 MMA Collector Buffer B Attribute"""

    B0 = 0
    B1 = 1
    B2 = 2
    B3 = 3

    def __str__(self):
        if self is Tcgen05MMACollectorBBuffer.B0:
            return "b0"
        if self is Tcgen05MMACollectorBBuffer.B1:
            return "b1"
        if self is Tcgen05MMACollectorBBuffer.B2:
            return "b2"
        if self is Tcgen05MMACollectorBBuffer.B3:
            return "b3"
        raise ValueError("Unknown Tcgen05MMACollectorBBuffer enum entry.")



@register_attribute_builder("Tcgen05MMACollectorBBuffer")
def _tcgen05mmacollectorbbuffer(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

class Tcgen05MMACollectorOp(IntEnum):
    """tcgen05.mma Collector Buffer Operation"""

    DISCARD = 0
    LASTUSE = 1
    FILL = 2
    USE = 3

    def __str__(self):
        if self is Tcgen05MMACollectorOp.DISCARD:
            return "discard"
        if self is Tcgen05MMACollectorOp.LASTUSE:
            return "lastuse"
        if self is Tcgen05MMACollectorOp.FILL:
            return "fill"
        if self is Tcgen05MMACollectorOp.USE:
            return "use"
        raise ValueError("Unknown Tcgen05MMACollectorOp enum entry.")



@register_attribute_builder("Tcgen05MMACollectorOp")
def _tcgen05mmacollectorop(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

class Tcgen05MMAKind(IntEnum):
    """tcgen05 MMA Supported Types"""

    F8F6F4 = 2
    I8 = 3
    F16 = 0
    TF32 = 1

    def __str__(self):
        if self is Tcgen05MMAKind.F8F6F4:
            return "f8f6f4"
        if self is Tcgen05MMAKind.I8:
            return "i8"
        if self is Tcgen05MMAKind.F16:
            return "f16"
        if self is Tcgen05MMAKind.TF32:
            return "tf32"
        raise ValueError("Unknown Tcgen05MMAKind enum entry.")



@register_attribute_builder("Tcgen05MMAKind")
def _tcgen05mmakind(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

class Tcgen05WaitKind(IntEnum):
    """NVVM Tcgen05 wait kind"""

    LOAD = 0
    STORE = 1

    def __str__(self):
        if self is Tcgen05WaitKind.LOAD:
            return "load"
        if self is Tcgen05WaitKind.STORE:
            return "store"
        raise ValueError("Unknown Tcgen05WaitKind enum entry.")



@register_attribute_builder("Tcgen05WaitKind")
def _tcgen05waitkind(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

class VoteSyncKind(IntEnum):
    """NVVM vote sync kind"""

    any = 0
    all = 1
    ballot = 2
    uni = 3

    def __str__(self):
        if self is VoteSyncKind.any:
            return "any"
        if self is VoteSyncKind.all:
            return "all"
        if self is VoteSyncKind.ballot:
            return "ballot"
        if self is VoteSyncKind.uni:
            return "uni"
        raise ValueError("Unknown VoteSyncKind enum entry.")



@register_attribute_builder("VoteSyncKind")
def _votesynckind(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

class WGMMAScaleIn(IntEnum):
    """WGMMA overflow options"""

    one = 1
    neg = auto()

    def __str__(self):
        if self is WGMMAScaleIn.one:
            return "one"
        if self is WGMMAScaleIn.neg:
            return "neg"
        raise ValueError("Unknown WGMMAScaleIn enum entry.")



@register_attribute_builder("WGMMAScaleIn")
def _wgmmascalein(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

class WGMMAScaleOut(IntEnum):
    """WGMMA input predicate"""

    zero = 0
    one = 1

    def __str__(self):
        if self is WGMMAScaleOut.zero:
            return "zero"
        if self is WGMMAScaleOut.one:
            return "one"
        raise ValueError("Unknown WGMMAScaleOut enum entry.")



@register_attribute_builder("WGMMAScaleOut")
def _wgmmascaleout(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

class WGMMATypes(IntEnum):
    """NVVM WGMMA types"""

    f16 = 0
    tf32 = 1
    u8 = 2
    s8 = 3
    b1 = 4
    bf16 = 5
    e4m3 = 6
    e5m2 = 7
    f32 = 8
    s32 = 9

    def __str__(self):
        if self is WGMMATypes.f16:
            return "f16"
        if self is WGMMATypes.tf32:
            return "tf32"
        if self is WGMMATypes.u8:
            return "u8"
        if self is WGMMATypes.s8:
            return "s8"
        if self is WGMMATypes.b1:
            return "b1"
        if self is WGMMATypes.bf16:
            return "bf16"
        if self is WGMMATypes.e4m3:
            return "e4m3"
        if self is WGMMATypes.e5m2:
            return "e5m2"
        if self is WGMMATypes.f32:
            return "f32"
        if self is WGMMATypes.s32:
            return "s32"
        raise ValueError("Unknown WGMMATypes enum entry.")



@register_attribute_builder("WGMMATypes")
def _wgmmatypes(x, context):
    return _ods_ir.IntegerAttr.get(_ods_ir.IntegerType.get_signless(32, context=context), int(x))

@register_attribute_builder("BarrierReductionAttr")
def _barrierreductionattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm.reduction<{str(x)}>', context=context)

@register_attribute_builder("BlockScaleFormatAttr")
def _blockscaleformatattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm.block_scale_format<{str(x)}>', context=context)

@register_attribute_builder("CTAGroupKindAttr")
def _ctagroupkindattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm.cta_group<{str(x)}>', context=context)

@register_attribute_builder("CacheEvictionPriorityAttr")
def _cacheevictionpriorityattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm<cache_eviction_priority {str(x)}>', context=context)

@register_attribute_builder("ClusterLaunchControlQueryTypeAttr")
def _clusterlaunchcontrolquerytypeattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm<cluster_launch_control_query_type {str(x)}>', context=context)

@register_attribute_builder("DotAccumulateTypeAttr")
def _dotaccumulatetypeattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm.dot_accumulate_type<{str(x)}>', context=context)

@register_attribute_builder("FPRoundingModeAttr")
def _fproundingmodeattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm.fp_rnd_mode<{str(x)}>', context=context)

@register_attribute_builder("GridDepActionAttr")
def _griddepactionattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm<grid_dep_action {str(x)}>', context=context)

@register_attribute_builder("LdStMatrixEltTypeAttr")
def _ldstmatrixelttypeattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm.ld_st_matrix_elt_type<{str(x)}>', context=context)

@register_attribute_builder("LoadCacheModifierAttr")
def _loadcachemodifierattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm<load_cache_modifier {str(x)}>', context=context)

@register_attribute_builder("MMAB1OpAttr")
def _mmab1opattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm.mma_b1op<{str(x)}>', context=context)

@register_attribute_builder("MMABlockScaleKindAttr")
def _mmablockscalekindattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm.block_scale_kind<{str(x)}>', context=context)

@register_attribute_builder("MMAFragAttr")
def _mmafragattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm.mma_frag<{str(x)}>', context=context)

@register_attribute_builder("MMAIntOverflowAttr")
def _mmaintoverflowattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm.mma_int_overflow<{str(x)}>', context=context)

@register_attribute_builder("MMAKindAttr")
def _mmakindattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm.mma_kind<{str(x)}>', context=context)

@register_attribute_builder("MMALayoutAttr")
def _mmalayoutattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm.mma_layout<{str(x)}>', context=context)

@register_attribute_builder("MMATypesAttr")
def _mmatypesattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm.mma_type<{str(x)}>', context=context)

@register_attribute_builder("MatchSyncKindAttr")
def _matchsynckindattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm<match_sync_kind {str(x)}>', context=context)

@register_attribute_builder("MemOrderKindAttr")
def _memorderkindattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm.mem_order<{str(x)}>', context=context)

@register_attribute_builder("MemScopeKindAttr")
def _memscopekindattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm.mem_scope<{str(x)}>', context=context)

@register_attribute_builder("NVVMMemorySpaceAttr")
def _nvvmmemoryspaceattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm.memory_space<{str(x)}>', context=context)

@register_attribute_builder("PermuteModeAttr")
def _permutemodeattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm.permute_mode<{str(x)}>', context=context)

@register_attribute_builder("PrefetchCacheLevelAttr")
def _prefetchcachelevelattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm<prefetch_cache_level {str(x)}>', context=context)

@register_attribute_builder("ProxyKindAttr")
def _proxykindattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm.proxy_kind<{str(x)}>', context=context)

@register_attribute_builder("ReduxKindAttr")
def _reduxkindattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm<redux_kind {str(x)}>', context=context)

@register_attribute_builder("SaturationModeAttr")
def _saturationmodeattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm.sat_mode<{str(x)}>', context=context)

@register_attribute_builder("ScaleVecSizeAttr")
def _scalevecsizeattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm.scale_vec_size<{str(x)}>', context=context)

@register_attribute_builder("SetMaxRegisterActionAttr")
def _setmaxregisteractionattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm<action {str(x)}>', context=context)

@register_attribute_builder("SharedSpaceAttr")
def _sharedspaceattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm.shared_space<{str(x)}>', context=context)

@register_attribute_builder("ShflKindAttr")
def _shflkindattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm<shfl_kind {str(x)}>', context=context)

@register_attribute_builder("TMALoadModeAttr")
def _tmaloadmodeattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm.tma_load_mode<{str(x)}>', context=context)

@register_attribute_builder("TMAReduxKindAttr")
def _tmareduxkindattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm.tma_redux_kind<{str(x)}>', context=context)

@register_attribute_builder("TMAStoreModeAttr")
def _tmastoremodeattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm.tma_store_mode<{str(x)}>', context=context)

@register_attribute_builder("Tcgen05CpMulticastAttr")
def _tcgen05cpmulticastattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm.tcgen05_cp_multicast<{str(x)}>', context=context)

@register_attribute_builder("Tcgen05CpShapeAttr")
def _tcgen05cpshapeattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm.tcgen05_cp_shape<{str(x)}>', context=context)

@register_attribute_builder("Tcgen05CpSrcFormatAttr")
def _tcgen05cpsrcformatattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm.tcgen05_cp_src_fmt<{str(x)}>', context=context)

@register_attribute_builder("Tcgen05FenceKindAttr")
def _tcgen05fencekindattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm.tcgen05_fence<{str(x)}>', context=context)

@register_attribute_builder("Tcgen05LdStShapeAttr")
def _tcgen05ldstshapeattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm.tcgen05_ldst_shape<{str(x)}>', context=context)

@register_attribute_builder("Tcgen05MMABlockScaleAttr")
def _tcgen05mmablockscaleattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm.tcgen05_mma_block_scale<{str(x)}>', context=context)

@register_attribute_builder("Tcgen05MMACollectorBBufferAttr")
def _tcgen05mmacollectorbbufferattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm.tcgen05_mma_collectorb<{str(x)}>', context=context)

@register_attribute_builder("Tcgen05MMACollectorOpAttr")
def _tcgen05mmacollectoropattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm.tcgen05_mma_collectorop<{str(x)}>', context=context)

@register_attribute_builder("Tcgen05MMAKindAttr")
def _tcgen05mmakindattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm.tcgen05_mma_kind<{str(x)}>', context=context)

@register_attribute_builder("Tcgen05WaitKindAttr")
def _tcgen05waitkindattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm.tcgen05_wait<{str(x)}>', context=context)

@register_attribute_builder("VoteSyncKindAttr")
def _votesynckindattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm<vote_sync_kind {str(x)}>', context=context)

@register_attribute_builder("WGMMAScaleInAttr")
def _wgmmascaleinattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm.wgmma_scale_in<{str(x)}>', context=context)

@register_attribute_builder("WGMMAScaleOutAttr")
def _wgmmascaleoutattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm.wgmma_scale_out<{str(x)}>', context=context)

@register_attribute_builder("WGMMATypesAttr")
def _wgmmatypesattr(x, context):
    return _ods_ir.Attribute.parse(f'#nvvm.wgmma_type<{str(x)}>', context=context)