singe-ptx 0.1.0-alpha.6

// auto-generated by `cargo xtask gen-ptx-instructions`
// do not edit manually

#[derive(Debug, Clone, Copy, PartialEq, Eq)]
pub enum InstructionOpcode {
    /// Absolute value.
    Abs,
    /// Queries the active threads within a warp.
    ActiveMask,
    /// Add two values.
    Add,
    /// Add two values with carry-out.
    AddCc,
    /// Add two values with carry-in and optional carry-out.
    Addc,
    /// Dynamically allocate memory on stack.
    Alloca,
    /// Bitwise AND.
    And,
    /// Apply the cache eviction priority to the specified address in the specified cache level.
    ApplyPriority,
    /// Atomic reduction operations for thread-to-thread communication.
    Atom,
    /// Barrier synchronization.
    Bar,
    /// Barrier synchronization for threads in a warp.
    BarWarpSync,
    /// Barrier synchronization.
    Barrier,
    /// Barrier synchronization within a cluster.
    BarrierCluster,
    /// Bit Field Extract.
    Bfe,
    /// Bit Field Insert.
    Bfi,
    /// Find most significant non-sign bit.
    Bfind,
    /// Bit Field Mask.
    Bmsk,
    /// Branch to a target and continue execution there.
    Bra,
    /// Bit reverse.
    Brev,
    /// Breakpoint.
    Brkpt,
    /// Branch to a label indexed from a list of potential branch targets.
    BrxIdx,
    /// Call a function, recording the return location.
    Call,
    /// Queries response of clusterlaunchcontrol.try_cancel operation.
    ClusterLaunchControlQueryCancel,
    /// Requests cancellation of cluster which is not launched yet.
    ClusterLaunchControlTryCancel,
    /// Count leading zeros.
    Clz,
    /// C/C++ style logical negation.
    Cnot,
    /// Copy sign of one input to another.
    Copysign,
    /// Find the cosine of a value.
    Cos,
    /// Makes the mbarrier object track all prior cp.async operations initiated by the executing thread.
    CpAsyncMbarrierArrive,
    /// Create a cache eviction policy for the specified cache level.
    CreatePolicy,
    /// Convert a value from one type to another.
    Cvt,
    /// Convert two integer values from one integer type to another and pack the results.
    CvtPack,
    /// Convert address from .const, Kernel Function Parameters ( .param ), .global, .local, or .shared state space to generic, or vice-versa. Take the generic address of a variable declared in .const, Kernel Function Parameters ( .param ), .global, .local, or .shared state space.
    Cvta,
    /// Discard the data at the specified address range and cache level.
    Discard,
    /// Divide one value by another.
    Div,
    /// Two-way dot product-accumulate.
    Dp2A,
    /// Four-way byte dot product-accumulate.
    Dp4A,
    /// Elect a leader thread from a set of threads.
    ElectSync,
    /// Find the base-2 exponential of a value.
    Ex2,
    /// Terminate a thread.
    Exit,
    /// Enforce an ordering of memory operations.
    Fence,
    /// Fused multiply-add.
    Fma,
    /// Find the n-th set bit
    Fns,
    /// Generate the CTA rank of the address.
    GetCtaRank,
    /// Control execution of dependent grids.
    GridDepControl,
    /// Query whether a generic address falls within a specified state space window.
    IsSpaceP,
    /// Query whether a register points to an opaque variable of a specified type.
    IsTypeP,
    /// Load a register variable from an addressable state space variable.
    Ld,
    /// Load a register variable from global state space via non-coherent cache.
    LdGlobalNc,
    /// Collectively load one or more matrices from shared memory for mma instruction
    LdMatrix,
    /// Load read-only data from an address that is common across threads in the warp.
    Ldu,
    /// Find the base-2 logarithm of a value.
    Lg2,
    /// Arbitrary logical operation on 3 inputs.
    Lop3,
    /// Multiply two values, optionally extract the high or low half of the intermediate result, and add a third value.
    Mad,
    /// Multiply two values, extract high or low half of result, and add a third value with carry-out.
    MadCc,
    /// Multiply two 24-bit integer values and add a third value.
    Mad24,
    /// Multiply two values, extract high or low half of result, and add a third value with carry-in and optional carry-out.
    Madc,
    /// Map the address of the shared variable in the target CTA.
    Mapa,
    /// Broadcast and compare a value across threads in warp.
    MatchSync,
    /// Find the maximum of two values.
    Max,
    /// Performs arrive-on operation on the mbarrier object .
    MbarrierArrive,
    /// Decrements the expected count of the mbarrier object and performs arrive-on operation .
    MbarrierArriveDrop,
    /// Performs complete-tx operation on the mbarrier object .
    MbarrierCompleteTx,
    /// Performs expect-tx operation on the mbarrier object .
    MbarrierExpectTx,
    /// Initialize the mbarrier object .
    MbarrierInit,
    /// Invalidates the mbarrier object .
    MbarrierInval,
    /// Query the pending arrival count from the opaque mbarrier state.
    MbarrierPendingCount,
    /// Checks whether the mbarrier object has completed the phase.
    MbarrierTestWait,
    /// Checks whether the mbarrier object has completed the phase.
    MbarrierTryWait,
    /// Enforce an ordering of memory operations.
    Membar,
    /// Find the minimum of two values.
    Min,
    /// Perform matrix multiply-and-accumulate operation
    Mma,
    /// Perform matrix multiply-and-accumulate operation with sparse matrix A
    MmaSp,
    /// Perform matrix multiply-and-accumulate operation with sparse matrix A
    MmaSpOrderedMetadata,
    /// Set a register variable with the value of a register variable or an immediate value. Take the non-generic address of a variable in global, local, or shared state space.
    Mov,
    /// Transpose a matrix in registers across the warp.
    MovMatrix,
    /// Multiply two values.
    Mul,
    /// Multiply two 24-bit integer values.
    Mul24,
    /// The multimem.* operations operate on multimem addresses and accesses all of the multiple memory locations which the multimem address points to.
    MultimemLdReduce,
    /// The multimem.* operations operate on multimem addresses and accesses all of the multiple memory locations which the multimem address points to.
    MultimemRed,
    /// The multimem.* operations operate on multimem addresses and accesses all of the multiple memory locations which the multimem address points to.
    MultimemSt,
    /// Suspend the thread for an approximate delay given in nanoseconds.
    Nanosleep,
    /// Arithmetic negate.
    Neg,
    /// Bitwise negation; one’s complement.
    Not,
    /// Biwise OR.
    Or,
    /// Trigger one or more Performance Monitor events.
    PmEvent,
    /// Population count.
    Popc,
    /// Prefetch line containing a generic address at a specified level of memory hierarchy, in specified state space.
    Prefetch,
    /// Prefetch line containing a generic address at a specified level of memory hierarchy, in specified state space.
    Prefetchu,
    /// Permute bytes from register pair.
    Prmt,
    /// Take the reciprocal of a value.
    Rcp,
    /// Compute a fast, gross approximation to the reciprocal of a value.
    RcpApproxFtzF64,
    /// Reduction operations on global and shared memory.
    Red,
    /// Asynchronous reduction operation.
    RedAsync,
    /// Perform reduction operation on the data from each predicated active thread in the thread group.
    ReduxSync,
    /// The remainder of integer division.
    Rem,
    /// Return from function to instruction after call.
    Ret,
    /// Take the reciprocal of the square root of a value.
    Rsqrt,
    /// Compute an approximation of the square root reciprocal of a value.
    RsqrtApproxFtzF64,
    /// Sum of absolute differences.
    Sad,
    /// Select between source operands, based on the value of the predicate source operand.
    Selp,
    /// Compare two numeric values with a relational operator, and optionally combine this result with a predicate value by applying a Boolean operator.
    Set,
    /// Hint to change the number of registers owned by the warp.
    SetMaxNReg,
    /// Compare two numeric values with a relational operator, and (optionally) combine this result with a predicate value by applying a Boolean operator.
    Setp,
    /// Funnel shift.
    Shf,
    /// Register data shuffle within threads of a warp.
    Shfl,
    /// Register data shuffle within threads of a warp.
    ShflSync,
    /// Shift bits left, zero-fill on right.
    Shl,
    /// Shift bits right, sign or zero-fill on left.
    Shr,
    /// Find the sine of a value.
    Sin,
    /// Select one source operand, based on the sign of the third operand.
    Slct,
    /// Take the square root of a value.
    Sqrt,
    /// Store data to an addressable state space variable.
    St,
    /// Asynchronous store operation.
    StAsync,
    /// Initializes a region of memory as specified by state space.
    StBulk,
    /// Update the stack pointer with a new value.
    StackRestore,
    /// Save the value of stack pointer into a register.
    StackSave,
    /// Collectively store one or more matrices to shared memory.
    StMatrix,
    /// Subtract one value from another.
    Sub,
    /// Subtract one value from another, with borrow-out.
    SubCc,
    /// Subtract one value from another, with borrow-in and optional borrow-out.
    Subc,
    /// Load from surface memory.
    Suld,
    /// Query a surface attribute.
    Suq,
    /// Reduce surface memory.
    Sured,
    /// Store to surface memory.
    Sust,
    /// Sign-extend or Zero-extend.
    Szext,
    /// Find the hyperbolic tangent of a value (in radians)
    Tanh,
    /// Dynamic Tensor Memory allocation management instructions
    Tcgen05Alloc,
    /// Makes the mbarrier object track the completion of all prior async-tcgen05 operations initiated by the executing thread.
    Tcgen05Commit,
    /// Initiates an asynchronous copy operation from shared memory to the Tensor Memory .
    Tcgen05Cp,
    /// Dynamic Tensor Memory allocation management instructions
    Tcgen05Dealloc,
    /// Specialized fence for the asynchronous tcgen05 operations.
    Tcgen05Fence,
    /// Asynchronous collective load from tensor memory into registers.
    Tcgen05Ld,
    /// Dynamic Tensor Memory allocation management instructions
    Tcgen05RelinquishAllocPermit,
    /// Asynchronously shift down the rows of the matrix in the Tensor Memory for a warp.
    Tcgen05Shift,
    /// Asynchronous collective store to tensor memory from registers.
    Tcgen05St,
    /// Waits for the completion of all prior asynchronous tcgen05.ld / tcgen05.st instructions.
    Tcgen05Wait,
    /// A fused copy and fence operation.
    TensorMapCpFenceProxy,
    /// Modifies the field of a tensor-map object.
    TensormapReplace,
    /// Test floating-point property.
    Testp,
    /// Perform a texture memory lookup.
    Tex,
    /// Perform a texture fetch of the 4-texel bilerp footprint.
    Tld4,
    /// Perform trap operation.
    Trap,
    /// Query texture and sampler attributes.
    Txq,
    /// Integer byte/half-word/word addition/subtraction.
    VAbsDiff,
    /// Integer dual half-word SIMD addition/subtraction.
    VAbsDiff2,
    /// Integer quad byte SIMD addition/subtraction.
    VAbsDiff4,
    /// Integer byte/half-word/word addition/subtraction.
    VAdd,
    /// Integer dual half-word SIMD addition/subtraction.
    VAdd2,
    /// Integer quad byte SIMD addition/subtraction.
    VAdd4,
    /// Integer dual half-word SIMD addition/subtraction.
    VAvrg2,
    /// Integer quad byte SIMD addition/subtraction.
    VAvrg4,
    /// Integer byte/half-word/word multiply-accumulate.
    VMad,
    /// Integer byte/half-word/word addition/subtraction.
    VMax,
    /// Integer dual half-word SIMD addition/subtraction.
    VMax2,
    /// Integer quad byte SIMD addition/subtraction.
    VMax4,
    /// Integer byte/half-word/word addition/subtraction.
    VMin,
    /// Integer dual half-word SIMD addition/subtraction.
    VMin2,
    /// Integer quad byte SIMD addition/subtraction.
    VMin4,
    /// Vote across thread group.
    Vote,
    /// Vote across thread group.
    VoteSync,
    /// Integer byte/half-word/word comparison.
    VSet,
    /// Integer dual half-word SIMD comparison.
    VSet2,
    /// Integer quad byte SIMD comparison.
    VSet4,
    /// Integer byte/half-word/word left/right shift.
    VShl,
    /// Integer byte/half-word/word left/right shift.
    VShr,
    /// Integer byte/half-word/word addition/subtraction.
    VSub,
    /// Integer dual half-word SIMD addition/subtraction.
    VSub2,
    /// Integer quad byte SIMD addition/subtraction.
    VSub4,
    /// Commits all prior uncommitted wgmma.mma_async operations into a wgmma-group .
    WgmmaCommitGroup,
    /// Enforce an ordering of register accesses between wgmma.mma_async and other operations.
    WgmmaFence,
    /// Perform matrix multiply-and-accumulate operation across warpgroup
    WgmmaMmaAsync,
    /// Perform matrix multiply-and-accumulate operation with sparse matrix A across warpgroup
    WgmmaMmaAsyncSp,
    /// Signal the completion of a preceding warpgroup operation.
    WgmmaWaitGroup,
    /// Collectively load a matrix from memory for WMMA
    WmmaLoad,
    /// Perform a single matrix multiply-and-accumulate operation across a warp
    WmmaMma,
    /// Collectively store a matrix into memory for WMMA
    WmmaStore,
    /// Bitwise exclusive-OR (inequality).
    Xor,
}
#[derive(Debug, Clone, PartialEq, Eq)]
pub struct InstructionSyntax {
    pub opcode: InstructionOpcode,
    pub name: &'static str,
    pub syntax: &'static str,
    pub min_operand_arity: usize,
    pub max_operand_arity: usize,
    pub required_modifiers: &'static [&'static str],
}
pub const PTX_91_INSTRUCTION_OPCODES: &[&str] = &[
    "abs",
    "activemask",
    "add",
    "add.cc",
    "addc",
    "alloca",
    "and",
    "applypriority",
    "atom",
    "bar",
    "bar.warp.sync",
    "barrier",
    "barrier.cluster",
    "bfe",
    "bfi",
    "bfind",
    "bmsk",
    "bra",
    "brev",
    "brkpt",
    "brx.idx",
    "call",
    "clusterlaunchcontrol.query_cancel",
    "clusterlaunchcontrol.try_cancel",
    "clz",
    "cnot",
    "copysign",
    "cos",
    "cp.async.mbarrier.arrive",
    "createpolicy",
    "cvt",
    "cvt.pack",
    "cvta",
    "discard",
    "div",
    "dp2a",
    "dp4a",
    "elect.sync",
    "ex2",
    "exit",
    "fence",
    "fma",
    "fns",
    "getctarank",
    "griddepcontrol",
    "isspacep",
    "istypep",
    "ld",
    "ld.global.nc",
    "ldmatrix",
    "ldu",
    "lg2",
    "lop3",
    "mad",
    "mad.cc",
    "mad24",
    "madc",
    "mapa",
    "match.sync",
    "max",
    "mbarrier.arrive",
    "mbarrier.arrive_drop",
    "mbarrier.complete_tx",
    "mbarrier.expect_tx",
    "mbarrier.init",
    "mbarrier.inval",
    "mbarrier.pending_count",
    "mbarrier.test_wait",
    "mbarrier.try_wait",
    "membar",
    "min",
    "mma",
    "mma.sp",
    "mma.sp::ordered_metadata",
    "mov",
    "movmatrix",
    "mul",
    "mul24",
    "multimem.ld_reduce",
    "multimem.red",
    "multimem.st",
    "nanosleep",
    "neg",
    "not",
    "or",
    "pmevent",
    "popc",
    "prefetch",
    "prefetchu",
    "prmt",
    "rcp",
    "rcp.approx.ftz.f64",
    "red",
    "red.async",
    "redux.sync",
    "rem",
    "ret",
    "rsqrt",
    "rsqrt.approx.ftz.f64",
    "sad",
    "selp",
    "set",
    "setmaxnreg",
    "setp",
    "shf",
    "shfl",
    "shfl.sync",
    "shl",
    "shr",
    "sin",
    "slct",
    "sqrt",
    "st",
    "st.async",
    "st.bulk",
    "stackrestore",
    "stacksave",
    "stmatrix",
    "sub",
    "sub.cc",
    "subc",
    "suld",
    "suq",
    "sured",
    "sust",
    "szext",
    "tanh",
    "tcgen05.alloc",
    "tcgen05.commit",
    "tcgen05.cp",
    "tcgen05.dealloc",
    "tcgen05.fence",
    "tcgen05.ld",
    "tcgen05.relinquish_alloc_permit",
    "tcgen05.shift",
    "tcgen05.st",
    "tcgen05.wait",
    "tensormap.cp_fenceproxy",
    "tensormap.replace",
    "testp",
    "tex",
    "tld4",
    "trap",
    "txq",
    "vabsdiff",
    "vabsdiff2",
    "vabsdiff4",
    "vadd",
    "vadd2",
    "vadd4",
    "vavrg2",
    "vavrg4",
    "vmad",
    "vmax",
    "vmax2",
    "vmax4",
    "vmin",
    "vmin2",
    "vmin4",
    "vote",
    "vote.sync",
    "vset",
    "vset2",
    "vset4",
    "vshl",
    "vshr",
    "vsub",
    "vsub2",
    "vsub4",
    "wgmma.commit_group",
    "wgmma.fence",
    "wgmma.mma_async",
    "wgmma.mma_async.sp",
    "wgmma.wait_group",
    "wmma.load",
    "wmma.mma",
    "wmma.store",
    "xor",
];
pub const PTX_91_INSTRUCTION_SYNTAX: &[InstructionSyntax] = &[
    InstructionSyntax {
        opcode: InstructionOpcode::Abs,
        name: "abs",
        syntax: "abs.bf16 d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".bf16"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Abs,
        name: "abs",
        syntax: "abs.bf16x2 d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".bf16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Abs,
        name: "abs",
        syntax: "abs.f64 d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".f64"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Abs,
        name: "abs",
        syntax: "abs.type d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Abs,
        name: "abs",
        syntax: "abs{.ftz}.f16 d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".f16"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Abs,
        name: "abs",
        syntax: "abs{.ftz}.f16x2 d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".f16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Abs,
        name: "abs",
        syntax: "abs{.ftz}.f32 d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::ActiveMask,
        name: "activemask",
        syntax: "activemask.b32 d;",
        min_operand_arity: 1usize,
        max_operand_arity: 1usize,
        required_modifiers: &[".b32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Add,
        name: "add",
        syntax: "add.type1 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Add,
        name: "add",
        syntax: "add{.rnd}.bf16 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".bf16"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Add,
        name: "add",
        syntax: "add{.rnd}.bf16x2 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".bf16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Add,
        name: "add",
        syntax: "add{.rnd}.f64 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".f64"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Add,
        name: "add",
        syntax: "add{.rnd}{.ftz}.f32x2 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".f32x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Add,
        name: "add",
        syntax: "add{.rnd}{.ftz}{.sat}.f16 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".f16"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Add,
        name: "add",
        syntax: "add{.rnd}{.ftz}{.sat}.f16x2 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".f16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Add,
        name: "add",
        syntax: "add{.rnd}{.ftz}{.sat}.f32 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Add,
        name: "add",
        syntax: "add{.rnd}{.sat}.f32.atype d, a, c;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Add,
        name: "add",
        syntax: "add{.sat}.type2 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::AddCc,
        name: "add.cc",
        syntax: "add.cc.type d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Addc,
        name: "addc",
        syntax: "addc{.cc}.type d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Alloca,
        name: "alloca",
        syntax: "alloca.type ptr, size{, immAlign};",
        min_operand_arity: 2usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::And,
        name: "and",
        syntax: "and.type d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::ApplyPriority,
        name: "applypriority",
        syntax: "applypriority{.global}.level::eviction_priority [a], size;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Atom,
        name: "atom",
        syntax: "atom{.sem}{.scope}{.space}.add.noftz{.level::cache_hint}.bf16 d, [a], b{, cache-policy};",
        min_operand_arity: 3usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".add", ".noftz", ".bf16"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Atom,
        name: "atom",
        syntax: "atom{.sem}{.scope}{.space}.add.noftz{.level::cache_hint}.bf16x2 d, [a], b{, cache-policy};",
        min_operand_arity: 3usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".add", ".noftz", ".bf16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Atom,
        name: "atom",
        syntax: "atom{.sem}{.scope}{.space}.add.noftz{.level::cache_hint}.f16 d, [a], b{, cache-policy};",
        min_operand_arity: 3usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".add", ".noftz", ".f16"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Atom,
        name: "atom",
        syntax: "atom{.sem}{.scope}{.space}.add.noftz{.level::cache_hint}.f16x2 d, [a], b{, cache-policy};",
        min_operand_arity: 3usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".add", ".noftz", ".f16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Atom,
        name: "atom",
        syntax: "atom{.sem}{.scope}{.space}.cas.b128 d, [a], b, c;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".cas", ".b128"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Atom,
        name: "atom",
        syntax: "atom{.sem}{.scope}{.space}.cas.b16 d, [a], b, c;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".cas", ".b16"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Atom,
        name: "atom",
        syntax: "atom{.sem}{.scope}{.space}.exch{.level::cache_hint}.b128 d, [a], b {, cache-policy};",
        min_operand_arity: 3usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".exch", ".b128"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Atom,
        name: "atom",
        syntax: "atom{.sem}{.scope}{.space}.op.type d, [a], b, c;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Atom,
        name: "atom",
        syntax: "atom{.sem}{.scope}{.space}.op{.level::cache_hint}.type d, [a], b{, cache-policy};",
        min_operand_arity: 3usize,
        max_operand_arity: 4usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Bar,
        name: "bar",
        syntax: "bar{.cta}.arrive a, b;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".arrive"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Bar,
        name: "bar",
        syntax: "bar{.cta}.red.op.pred p, a{, b}, {!}c;",
        min_operand_arity: 2usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".red", ".pred"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Bar,
        name: "bar",
        syntax: "bar{.cta}.red.popc.u32 d, a{, b}, {!}c;",
        min_operand_arity: 2usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".red", ".popc", ".u32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Bar,
        name: "bar",
        syntax: "bar{.cta}.sync a{, b};",
        min_operand_arity: 1usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".sync"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::BarWarpSync,
        name: "bar.warp.sync",
        syntax: "bar.warp.sync membermask;",
        min_operand_arity: 1usize,
        max_operand_arity: 1usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Barrier,
        name: "barrier",
        syntax: "barrier{.cta}.arrive{.aligned} a, b;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".arrive"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Barrier,
        name: "barrier",
        syntax: "barrier{.cta}.red.op{.aligned}.pred p, a{, b}, {!}c;",
        min_operand_arity: 2usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".red", ".pred"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Barrier,
        name: "barrier",
        syntax: "barrier{.cta}.red.popc{.aligned}.u32 d, a{, b}, {!}c;",
        min_operand_arity: 2usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".red", ".popc", ".u32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Barrier,
        name: "barrier",
        syntax: "barrier{.cta}.sync{.aligned} a{, b};",
        min_operand_arity: 1usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".sync"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::BarrierCluster,
        name: "barrier.cluster",
        syntax: "barrier.cluster.arrive{.sem}{.aligned};",
        min_operand_arity: 0usize,
        max_operand_arity: 0usize,
        required_modifiers: &[".arrive"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::BarrierCluster,
        name: "barrier.cluster",
        syntax: "barrier.cluster.wait{.acquire}{.aligned};",
        min_operand_arity: 0usize,
        max_operand_arity: 0usize,
        required_modifiers: &[".wait"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Bfe,
        name: "bfe",
        syntax: "bfe.type d, a, b, c;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Bfi,
        name: "bfi",
        syntax: "bfi.type f, a, b, c, d;",
        min_operand_arity: 5usize,
        max_operand_arity: 5usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Bfind,
        name: "bfind",
        syntax: "bfind.shiftamt.type d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".shiftamt"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Bfind,
        name: "bfind",
        syntax: "bfind.type d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Bmsk,
        name: "bmsk",
        syntax: "bmsk.mode.b32 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".b32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Brev,
        name: "brev",
        syntax: "brev.type d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Brkpt,
        name: "brkpt",
        syntax: "brkpt;",
        min_operand_arity: 0usize,
        max_operand_arity: 0usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::BrxIdx,
        name: "brx.idx",
        syntax: "brx.idx{.uni} index, tlist;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Call,
        name: "call",
        syntax: "call{.uni} (ret-param), fptr, (param-list), flist;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Call,
        name: "call",
        syntax: "call{.uni} (ret-param), fptr, (param-list), fproto;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Call,
        name: "call",
        syntax: "call{.uni} (ret-param), func, (param-list);",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Call,
        name: "call",
        syntax: "call{.uni} fptr, (param-list), flist;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Call,
        name: "call",
        syntax: "call{.uni} fptr, (param-list), fproto;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Call,
        name: "call",
        syntax: "call{.uni} fptr, flist;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Call,
        name: "call",
        syntax: "call{.uni} fptr, fproto;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Call,
        name: "call",
        syntax: "call{.uni} func, (param-list);",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Call,
        name: "call",
        syntax: "call{.uni} func;",
        min_operand_arity: 1usize,
        max_operand_arity: 1usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::ClusterLaunchControlQueryCancel,
        name: "clusterlaunchcontrol.query_cancel",
        syntax: "clusterlaunchcontrol.query_cancel.get_first_ctaid.v4.b32.b128 {xdim, ydim, zdim, _}, try_cancel_response;",
        min_operand_arity: 4usize,
        max_operand_arity: 5usize,
        required_modifiers: &[".get_first_ctaid", ".v4", ".b32", ".b128"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::ClusterLaunchControlQueryCancel,
        name: "clusterlaunchcontrol.query_cancel",
        syntax: "clusterlaunchcontrol.query_cancel.get_first_ctaid{::dimension}.b32.b128 reg, try_cancel_response;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".get_first_ctaid", ".b32", ".b128"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::ClusterLaunchControlQueryCancel,
        name: "clusterlaunchcontrol.query_cancel",
        syntax: "clusterlaunchcontrol.query_cancel.is_canceled.pred.b128 pred, try_cancel_response;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".is_canceled", ".pred", ".b128"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::ClusterLaunchControlTryCancel,
        name: "clusterlaunchcontrol.try_cancel",
        syntax: "clusterlaunchcontrol.try_cancel.async{.space}.completion_mechanism{.multicast::cluster::all}.b128 [addr], [mbar];",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".async", ".b128"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Clz,
        name: "clz",
        syntax: "clz.type d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Cnot,
        name: "cnot",
        syntax: "cnot.type d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Copysign,
        name: "copysign",
        syntax: "copysign.type d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Cos,
        name: "cos",
        syntax: "cos.approx{.ftz}.f32 d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".approx", ".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::CpAsyncMbarrierArrive,
        name: "cp.async.mbarrier.arrive",
        syntax: "cp.async.mbarrier.arrive{.noinc}{.shared{::cta}}.b64 [addr];",
        min_operand_arity: 1usize,
        max_operand_arity: 1usize,
        required_modifiers: &[".b64"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::CreatePolicy,
        name: "createpolicy",
        syntax: "createpolicy.cvt.L2.b64 cache-policy, access-property;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".cvt", ".L2", ".b64"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::CreatePolicy,
        name: "createpolicy",
        syntax: "createpolicy.fractional.level::primary_priority{.level::secondary_priority}.b64 cache-policy{, fraction};",
        min_operand_arity: 1usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".fractional", ".b64"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::CreatePolicy,
        name: "createpolicy",
        syntax: "createpolicy.range{.global}.level::primary_priority{.level::secondary_priority}.b64 cache-policy, [a], primary-size, total-size;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".range", ".b64"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Cvt,
        name: "cvt",
        syntax: "cvt.frnd2{.relu}{.satfinite}.bf16.f32 d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".bf16", ".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Cvt,
        name: "cvt",
        syntax: "cvt.frnd2{.relu}{.satfinite}.bf16x2.f32 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".bf16x2", ".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Cvt,
        name: "cvt",
        syntax: "cvt.frnd2{.relu}{.satfinite}.f16.f32 d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".f16", ".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Cvt,
        name: "cvt",
        syntax: "cvt.frnd2{.relu}{.satfinite}.f16x2.f32 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".f16x2", ".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Cvt,
        name: "cvt",
        syntax: "cvt.frnd2{.satfinite}{.relu}.tf32.f32 d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".tf32", ".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Cvt,
        name: "cvt",
        syntax: "cvt.frnd3{.satfinite}.ue8m0x2.bf16x2 d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".ue8m0x2", ".bf16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Cvt,
        name: "cvt",
        syntax: "cvt.frnd3{.satfinite}.ue8m0x2.f32 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".ue8m0x2", ".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Cvt,
        name: "cvt",
        syntax: "cvt.rn.bf16x2.ue8m0x2 d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".rn", ".bf16x2", ".ue8m0x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Cvt,
        name: "cvt",
        syntax: "cvt.rn.satfinite{.relu}.f4x2type.f32 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".rn", ".satfinite", ".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Cvt,
        name: "cvt",
        syntax: "cvt.rn.satfinite{.relu}.f4x2type.fp16x2type d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".rn", ".satfinite"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Cvt,
        name: "cvt",
        syntax: "cvt.rn.satfinite{.relu}.f6x2type.f32 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".rn", ".satfinite", ".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Cvt,
        name: "cvt",
        syntax: "cvt.rn.satfinite{.relu}.f6x2type.fp16x2type d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".rn", ".satfinite"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Cvt,
        name: "cvt",
        syntax: "cvt.rn.satfinite{.relu}.f8x2type.f32 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".rn", ".satfinite", ".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Cvt,
        name: "cvt",
        syntax: "cvt.rn.satfinite{.relu}.f8x2type.fp16x2 d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".rn", ".satfinite", ".fp16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Cvt,
        name: "cvt",
        syntax: "cvt.rn.satfinite{.relu}{.scaled::n2::ue8m0}.s2f6x2.bf16x2 d, a{, scale-factor};",
        min_operand_arity: 2usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".rn", ".satfinite", ".s2f6x2", ".bf16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Cvt,
        name: "cvt",
        syntax: "cvt.rn.satfinite{.relu}{.scaled::n2::ue8m0}.s2f6x2.f32 d, a, b{, scale-factor};",
        min_operand_arity: 3usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".rn", ".satfinite", ".s2f6x2", ".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Cvt,
        name: "cvt",
        syntax: "cvt.rna{.satfinite}.tf32.f32 d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".rna", ".tf32", ".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Cvt,
        name: "cvt",
        syntax: "cvt.rn{.relu}.f16x2.f4x2type d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".rn", ".f16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Cvt,
        name: "cvt",
        syntax: "cvt.rn{.relu}.f16x2.f6x2type d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".rn", ".f16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Cvt,
        name: "cvt",
        syntax: "cvt.rn{.relu}.f16x2.f8x2type d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".rn", ".f16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Cvt,
        name: "cvt",
        syntax: "cvt.rn{.relu}{.satfinite}{.scaled::n2::ue8m0}.bf16x2.f4x2type d, a{, scale-factor};",
        min_operand_arity: 2usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".rn", ".bf16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Cvt,
        name: "cvt",
        syntax: "cvt.rn{.relu}{.satfinite}{.scaled::n2::ue8m0}.bf16x2.f6x2type d, a{, scale-factor};",
        min_operand_arity: 2usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".rn", ".bf16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Cvt,
        name: "cvt",
        syntax: "cvt.rn{.relu}{.satfinite}{.scaled::n2::ue8m0}.bf16x2.f8x2type d, a{, scale-factor};",
        min_operand_arity: 2usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".rn", ".bf16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Cvt,
        name: "cvt",
        syntax: "cvt.rn{.satfinite}{.relu}{.scaled::n2::ue8m0}.bf16x2.s2f6x2 d, a{, scale-factor};",
        min_operand_arity: 2usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".rn", ".bf16x2", ".s2f6x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Cvt,
        name: "cvt",
        syntax: "cvt.rs{.relu}.satfinite.f4x4type.f32 d, {a, b, e, f}, rbits;",
        min_operand_arity: 5usize,
        max_operand_arity: 6usize,
        required_modifiers: &[".rs", ".satfinite", ".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Cvt,
        name: "cvt",
        syntax: "cvt.rs{.relu}.satfinite.f6x4type.f32 d, {a, b, e, f}, rbits;",
        min_operand_arity: 5usize,
        max_operand_arity: 6usize,
        required_modifiers: &[".rs", ".satfinite", ".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Cvt,
        name: "cvt",
        syntax: "cvt.rs{.relu}.satfinite.f8x4type.f32 d, {a, b, e, f}, rbits;",
        min_operand_arity: 5usize,
        max_operand_arity: 6usize,
        required_modifiers: &[".rs", ".satfinite", ".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Cvt,
        name: "cvt",
        syntax: "cvt.rs{.relu}{.satfinite}.bf16x2.f32 d, a, b, rbits;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".rs", ".bf16x2", ".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Cvt,
        name: "cvt",
        syntax: "cvt.rs{.relu}{.satfinite}.f16x2.f32 d, a, b, rbits;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".rs", ".f16x2", ".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Cvt,
        name: "cvt",
        syntax: "cvt{.irnd}{.ftz}{.sat}.dtype.atype d, a; cvt{.frnd}{.ftz}{.sat}.dtype.atype d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::CvtPack,
        name: "cvt.pack",
        syntax: "cvt.pack.sat.convertType.abType d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".sat"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::CvtPack,
        name: "cvt.pack",
        syntax: "cvt.pack.sat.convertType.abType.cType d, a, b, c;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".sat"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Cvta,
        name: "cvta",
        syntax: "cvta.space.size p, a; cvta.space.size p, var; cvta.space.size p, var+imm;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Cvta,
        name: "cvta",
        syntax: "cvta.to.space.size p, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".to"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Discard,
        name: "discard",
        syntax: "discard{.global}.level [a], size;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Div,
        name: "div",
        syntax: "div.approx{.ftz}.f32 d, a, b; div.full{.ftz}.f32 d, a, b; div.rnd{.ftz}.f32 d, a, b; div.rnd.f64 d, a, b;",
        min_operand_arity: 7usize,
        max_operand_arity: 9usize,
        required_modifiers: &[".approx", ".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Div,
        name: "div",
        syntax: "div.type d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Dp2A,
        name: "dp2a",
        syntax: "dp2a.mode.atype.btype d, a, b, c;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Dp4A,
        name: "dp4a",
        syntax: "dp4a.atype.btype d, a, b, c;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::ElectSync,
        name: "elect.sync",
        syntax: "elect.sync d|p, membermask;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Ex2,
        name: "ex2",
        syntax: "ex2.approx.atype d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".approx"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Ex2,
        name: "ex2",
        syntax: "ex2.approx.ftz.btype d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".approx", ".ftz"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Ex2,
        name: "ex2",
        syntax: "ex2.approx{.ftz}.f32 d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".approx", ".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Exit,
        name: "exit",
        syntax: "exit;",
        min_operand_arity: 0usize,
        max_operand_arity: 0usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Fence,
        name: "fence",
        syntax: "fence.acquire.sync_restrict::shared::cluster.cluster;",
        min_operand_arity: 0usize,
        max_operand_arity: 0usize,
        required_modifiers: &[".acquire", ".sync_restrict::shared::cluster", ".cluster"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Fence,
        name: "fence",
        syntax: "fence.op_restrict.release.cluster;",
        min_operand_arity: 0usize,
        max_operand_arity: 0usize,
        required_modifiers: &[".release", ".cluster"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Fence,
        name: "fence",
        syntax: "fence.proxy.async::generic.acquire.sync_restrict::shared::cluster.cluster;",
        min_operand_arity: 0usize,
        max_operand_arity: 0usize,
        required_modifiers: &[
            ".proxy",
            ".async::generic",
            ".acquire",
            ".sync_restrict::shared::cluster",
            ".cluster",
        ],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Fence,
        name: "fence",
        syntax: "fence.proxy.async::generic.release.sync_restrict::shared::cta.cluster;",
        min_operand_arity: 0usize,
        max_operand_arity: 0usize,
        required_modifiers: &[
            ".proxy",
            ".async::generic",
            ".release",
            ".sync_restrict::shared::cta",
            ".cluster",
        ],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Fence,
        name: "fence",
        syntax: "fence.proxy.proxykind;",
        min_operand_arity: 0usize,
        max_operand_arity: 0usize,
        required_modifiers: &[".proxy"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Fence,
        name: "fence",
        syntax: "fence.proxy.to_proxykind::from_proxykind.acquire.scope [addr], size;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".proxy", ".acquire"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Fence,
        name: "fence",
        syntax: "fence.proxy.to_proxykind::from_proxykind.release.scope;",
        min_operand_arity: 0usize,
        max_operand_arity: 0usize,
        required_modifiers: &[".proxy", ".release"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Fence,
        name: "fence",
        syntax: "fence.release.sync_restrict::shared::cta.cluster;",
        min_operand_arity: 0usize,
        max_operand_arity: 0usize,
        required_modifiers: &[".release", ".sync_restrict::shared::cta", ".cluster"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Fence,
        name: "fence",
        syntax: "fence{.sem}.scope;",
        min_operand_arity: 0usize,
        max_operand_arity: 0usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Fma,
        name: "fma",
        syntax: "fma.rnd.f64 d, a, b, c;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".f64"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Fma,
        name: "fma",
        syntax: "fma.rnd.oob.{relu}.type d, a, b, c;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".oob", ".", ".type"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Fma,
        name: "fma",
        syntax: "fma.rnd{.ftz}.f32x2 d, a, b, c;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".f32x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Fma,
        name: "fma",
        syntax: "fma.rnd{.ftz}.relu.f16 d, a, b, c;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".relu", ".f16"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Fma,
        name: "fma",
        syntax: "fma.rnd{.ftz}.relu.f16x2 d, a, b, c;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".relu", ".f16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Fma,
        name: "fma",
        syntax: "fma.rnd{.ftz}{.sat}.f16 d, a, b, c;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".f16"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Fma,
        name: "fma",
        syntax: "fma.rnd{.ftz}{.sat}.f16x2 d, a, b, c;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".f16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Fma,
        name: "fma",
        syntax: "fma.rnd{.ftz}{.sat}.f32 d, a, b, c;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Fma,
        name: "fma",
        syntax: "fma.rnd{.relu}.bf16 d, a, b, c;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".bf16"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Fma,
        name: "fma",
        syntax: "fma.rnd{.relu}.bf16x2 d, a, b, c;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".bf16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Fma,
        name: "fma",
        syntax: "fma.rnd{.sat}.f32.abtype d, a, b, c;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Fns,
        name: "fns",
        syntax: "fns.b32 d, mask, base, offset;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".b32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::GetCtaRank,
        name: "getctarank",
        syntax: "getctarank.shared::cluster.type d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".shared::cluster"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::GetCtaRank,
        name: "getctarank",
        syntax: "getctarank.shared::cluster.type d, var + imm;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".shared::cluster"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::GetCtaRank,
        name: "getctarank",
        syntax: "getctarank.shared::cluster.type d, var;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".shared::cluster"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::GetCtaRank,
        name: "getctarank",
        syntax: "getctarank.type d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::GetCtaRank,
        name: "getctarank",
        syntax: "getctarank{.space}.type d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::GridDepControl,
        name: "griddepcontrol",
        syntax: "griddepcontrol.action;",
        min_operand_arity: 0usize,
        max_operand_arity: 0usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::IsSpaceP,
        name: "isspacep",
        syntax: "isspacep.space p, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::IsTypeP,
        name: "istypep",
        syntax: "istypep.type p, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Ld,
        name: "ld",
        syntax: "ld.acquire.scope{.ss}{.level1::eviction_priority}{.level2::eviction_priority}{.level::cache_hint}{.level::prefetch_size}{.vec}.type d, [a]{, cache-policy};",
        min_operand_arity: 2usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".acquire"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Ld,
        name: "ld",
        syntax: "ld.mmio.relaxed.sys{.global}.type d, [a];",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".mmio", ".relaxed", ".sys"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Ld,
        name: "ld",
        syntax: "ld.relaxed.scope{.ss}{.level1::eviction_priority}{.level2::eviction_priority}{.level::cache_hint}{.level::prefetch_size}{.vec}.type d, [a]{, cache-policy};",
        min_operand_arity: 2usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".relaxed"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Ld,
        name: "ld",
        syntax: "ld.volatile{.ss}{.level::prefetch_size}{.vec}.type d, [a];",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".volatile"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Ld,
        name: "ld",
        syntax: "ld{.weak}{.ss}{.cop}{.level::cache_hint}{.level::prefetch_size}{.vec}.type d, [a]{.unified}{, cache-policy};",
        min_operand_arity: 2usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Ld,
        name: "ld",
        syntax: "ld{.weak}{.ss}{.level1::eviction_priority}{.level2::eviction_priority}{.level::cache_hint}{.level::prefetch_size}{.vec}.type d, [a]{.unified}{, cache-policy};",
        min_operand_arity: 2usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::LdGlobalNc,
        name: "ld.global.nc",
        syntax: "ld.global.nc{.level1::eviction_priority}{.level2::eviction_priority}{.level::cache_hint}{.level::prefetch_size}.type d, [a]{, cache-policy};",
        min_operand_arity: 2usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::LdGlobalNc,
        name: "ld.global.nc",
        syntax: "ld.global.nc{.level1::eviction_priority}{.level2::eviction_priority}{.level::cache_hint}{.level::prefetch_size}.vec.type d, [a]{, cache-policy};",
        min_operand_arity: 2usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::LdMatrix,
        name: "ldmatrix",
        syntax: "ldmatrix.sync.aligned.m16n16.num.trans{.ss}.dst_fmt.src_fmt r, [p];",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".sync", ".aligned", ".m16n16", ".trans"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::LdMatrix,
        name: "ldmatrix",
        syntax: "ldmatrix.sync.aligned.m8n16.num{.ss}.dst_fmt.src_fmt r, [p];",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".sync", ".aligned", ".m8n16"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::LdMatrix,
        name: "ldmatrix",
        syntax: "ldmatrix.sync.aligned.shape.num{.trans}{.ss}.type r, [p];",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".sync", ".aligned"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Ldu,
        name: "ldu",
        syntax: "ldu{.ss}.type d, [a]; ldu{.ss}.vec.type d, [a];",
        min_operand_arity: 2usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Lg2,
        name: "lg2",
        syntax: "lg2.approx{.ftz}.f32 d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".approx", ".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Lop3,
        name: "lop3",
        syntax: "lop3.BoolOp.b32 d|p, a, b, c, immLut, q;",
        min_operand_arity: 6usize,
        max_operand_arity: 6usize,
        required_modifiers: &[".b32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Lop3,
        name: "lop3",
        syntax: "lop3.b32 d, a, b, c, immLut;",
        min_operand_arity: 5usize,
        max_operand_arity: 5usize,
        required_modifiers: &[".b32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Mad,
        name: "mad",
        syntax: "mad.hi.sat.s32 d, a, b, c;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".hi", ".sat", ".s32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Mad,
        name: "mad",
        syntax: "mad.mode.type d, a, b, c;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Mad,
        name: "mad",
        syntax: "mad{.ftz}{.sat}.f32 d, a, b, c; mad.rnd{.ftz}{.sat}.f32 d, a, b, c; mad.rnd.f64 d, a, b, c;",
        min_operand_arity: 9usize,
        max_operand_arity: 10usize,
        required_modifiers: &[".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Mad24,
        name: "mad24",
        syntax: "mad24.hi.sat.s32 d, a, b, c;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".hi", ".sat", ".s32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Mad24,
        name: "mad24",
        syntax: "mad24.mode.type d, a, b, c;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Madc,
        name: "madc",
        syntax: "madc{.hi,.lo}{.cc}.type d, a, b, c;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Mapa,
        name: "mapa",
        syntax: "mapa.shared::cluster.type d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".shared::cluster"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Mapa,
        name: "mapa",
        syntax: "mapa.shared::cluster.type d, sh + imm, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".shared::cluster"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Mapa,
        name: "mapa",
        syntax: "mapa.shared::cluster.type d, sh, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".shared::cluster"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Mapa,
        name: "mapa",
        syntax: "mapa.type d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Mapa,
        name: "mapa",
        syntax: "mapa{.space}.type d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Max,
        name: "max",
        syntax: "max.f64 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".f64"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Max,
        name: "max",
        syntax: "max.type1 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Max,
        name: "max",
        syntax: "max{.NaN}{.xorsign.abs}.bf16 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".bf16"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Max,
        name: "max",
        syntax: "max{.NaN}{.xorsign.abs}.bf16x2 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".bf16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Max,
        name: "max",
        syntax: "max{.ftz}{.NaN}{.abs}.f32 d, a, b, c;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Max,
        name: "max",
        syntax: "max{.ftz}{.NaN}{.xorsign.abs}.f16 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".f16"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Max,
        name: "max",
        syntax: "max{.ftz}{.NaN}{.xorsign.abs}.f16x2 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".f16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Max,
        name: "max",
        syntax: "max{.ftz}{.NaN}{.xorsign.abs}.f32 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Max,
        name: "max",
        syntax: "max{.relu}.type2 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::MbarrierArrive,
        name: "mbarrier.arrive",
        syntax: "mbarrier.arrive.expect_tx{.sem.scope}{.shared::cluster}.b64 _, [addr], txCount;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".expect_tx", ".b64"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::MbarrierArrive,
        name: "mbarrier.arrive",
        syntax: "mbarrier.arrive.noComplete{.release.cta}{.shared{::cta}}.b64 state, [addr], count;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".noComplete", ".b64"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::MbarrierArrive,
        name: "mbarrier.arrive",
        syntax: "mbarrier.arrive{.sem.scope}{.shared::cluster}.b64 _, [addr] {,count} mbarrier.arrive.expect_tx{.sem.scope}{.shared{::cta}}.b64 state, [addr], txCount;",
        min_operand_arity: 3usize,
        max_operand_arity: 5usize,
        required_modifiers: &[".b64"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::MbarrierArrive,
        name: "mbarrier.arrive",
        syntax: "mbarrier.arrive{.sem.scope}{.shared{::cta}}.b64 state, [addr]{, count};",
        min_operand_arity: 2usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".b64"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::MbarrierArriveDrop,
        name: "mbarrier.arrive_drop",
        syntax: "mbarrier.arrive_drop.expect_tx{.sem.scope}{.shared::cluster}.b64 _, [addr], tx_count;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".expect_tx", ".b64"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::MbarrierArriveDrop,
        name: "mbarrier.arrive_drop",
        syntax: "mbarrier.arrive_drop.expect_tx{.sem.scope}{.shared{::cta}}.b64 state, [addr], tx_count;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".expect_tx", ".b64"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::MbarrierArriveDrop,
        name: "mbarrier.arrive_drop",
        syntax: "mbarrier.arrive_drop.noComplete{.release.cta}{.shared{::cta}}.b64 state, [addr], count;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".noComplete", ".b64"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::MbarrierArriveDrop,
        name: "mbarrier.arrive_drop",
        syntax: "mbarrier.arrive_drop{.sem.scope}{.shared::cluster}.b64 _, [addr] {, count};",
        min_operand_arity: 2usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".b64"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::MbarrierArriveDrop,
        name: "mbarrier.arrive_drop",
        syntax: "mbarrier.arrive_drop{.sem.scope}{.shared{::cta}}.b64 state, [addr] {, count};",
        min_operand_arity: 2usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".b64"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::MbarrierCompleteTx,
        name: "mbarrier.complete_tx",
        syntax: "mbarrier.complete_tx{.sem.scope}{.space}.b64 [addr], txCount;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".b64"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::MbarrierExpectTx,
        name: "mbarrier.expect_tx",
        syntax: "mbarrier.expect_tx{.sem.scope}{.space}.b64 [addr], txCount;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".b64"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::MbarrierInit,
        name: "mbarrier.init",
        syntax: "mbarrier.init{.shared{::cta}}.b64 [addr], count;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".b64"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::MbarrierInval,
        name: "mbarrier.inval",
        syntax: "mbarrier.inval{.shared{::cta}}.b64 [addr];",
        min_operand_arity: 1usize,
        max_operand_arity: 1usize,
        required_modifiers: &[".b64"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::MbarrierPendingCount,
        name: "mbarrier.pending_count",
        syntax: "mbarrier.pending_count.b64 count, state;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".b64"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::MbarrierTestWait,
        name: "mbarrier.test_wait",
        syntax: "mbarrier.test_wait.parity{.sem.scope}{.shared{::cta}}.b64 waitComplete, [addr], phaseParity;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".parity", ".b64"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::MbarrierTestWait,
        name: "mbarrier.test_wait",
        syntax: "mbarrier.test_wait{.sem.scope}{.shared{::cta}}.b64 waitComplete, [addr], state;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".b64"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::MbarrierTryWait,
        name: "mbarrier.try_wait",
        syntax: "mbarrier.try_wait.parity{.sem.scope}{.shared{::cta}}.b64 waitComplete, [addr], phaseParity {, timeHint};",
        min_operand_arity: 3usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".parity", ".b64"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::MbarrierTryWait,
        name: "mbarrier.try_wait",
        syntax: "mbarrier.try_wait{.sem.scope}{.shared{::cta}}.b64 waitComplete, [addr], state {, timeHint};",
        min_operand_arity: 3usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".b64"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Membar,
        name: "membar",
        syntax: "membar.level;",
        min_operand_arity: 0usize,
        max_operand_arity: 0usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Membar,
        name: "membar",
        syntax: "membar.proxy.proxykind;",
        min_operand_arity: 0usize,
        max_operand_arity: 0usize,
        required_modifiers: &[".proxy"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Min,
        name: "min",
        syntax: "min.f64 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".f64"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Min,
        name: "min",
        syntax: "min.type1 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Min,
        name: "min",
        syntax: "min{.NaN}{.xorsign.abs}.bf16 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".bf16"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Min,
        name: "min",
        syntax: "min{.NaN}{.xorsign.abs}.bf16x2 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".bf16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Min,
        name: "min",
        syntax: "min{.ftz}{.NaN}{.abs}.f32 d, a, b, c;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Min,
        name: "min",
        syntax: "min{.ftz}{.NaN}{.xorsign.abs}.f16 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".f16"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Min,
        name: "min",
        syntax: "min{.ftz}{.NaN}{.xorsign.abs}.f16x2 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".f16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Min,
        name: "min",
        syntax: "min{.ftz}{.NaN}{.xorsign.abs}.f32 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Min,
        name: "min",
        syntax: "min{.relu}.type2 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Mma,
        name: "mma",
        syntax: "mma.sync.aligned.m16n8k16.row.col.dtype.f16.f16.ctype d, a, b, c;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[
            ".sync",
            ".aligned",
            ".m16n8k16",
            ".row",
            ".col",
            ".f16",
            ".f16",
        ],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Mma,
        name: "mma",
        syntax: "mma.sync.aligned.m16n8k8.row.col.dtype.f16.f16.ctype d, a, b, c;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[
            ".sync", ".aligned", ".m16n8k8", ".row", ".col", ".f16", ".f16",
        ],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Mma,
        name: "mma",
        syntax: "mma.sync.aligned.m8n8k4.alayout.blayout.dtype.f16.f16.ctype d, a, b, c;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".sync", ".aligned", ".m8n8k4", ".f16", ".f16"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Mov,
        name: "mov",
        syntax: "mov.type d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Mov,
        name: "mov",
        syntax: "mov.type d, avar; mov.type d, avar+imm; mov.u32 d, fname; mov.u64 d, fname; mov.u32 d, kernel; mov.u64 d, kernel;",
        min_operand_arity: 7usize,
        max_operand_arity: 7usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Mov,
        name: "mov",
        syntax: "mov.type d, sreg;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::MovMatrix,
        name: "movmatrix",
        syntax: "movmatrix.sync.aligned.shape.trans.type d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".sync", ".aligned", ".trans"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Mul,
        name: "mul",
        syntax: "mul.mode.type d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Mul,
        name: "mul",
        syntax: "mul{.rnd}.bf16 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".bf16"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Mul,
        name: "mul",
        syntax: "mul{.rnd}.bf16x2 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".bf16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Mul,
        name: "mul",
        syntax: "mul{.rnd}.f64 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".f64"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Mul,
        name: "mul",
        syntax: "mul{.rnd}{.ftz}.f32x2 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".f32x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Mul,
        name: "mul",
        syntax: "mul{.rnd}{.ftz}{.sat}.f16 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".f16"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Mul,
        name: "mul",
        syntax: "mul{.rnd}{.ftz}{.sat}.f16x2 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".f16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Mul,
        name: "mul",
        syntax: "mul{.rnd}{.ftz}{.sat}.f32 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Mul24,
        name: "mul24",
        syntax: "mul24.mode.type d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::MultimemLdReduce,
        name: "multimem.ld_reduce",
        syntax: "multimem.ld_reduce.weak{.ss}.op.type d, [a];",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".weak"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::MultimemLdReduce,
        name: "multimem.ld_reduce",
        syntax: "multimem.ld_reduce.weak{.ss}.op{.acc_prec}{.vec}.type d, [a];",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".weak"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::MultimemLdReduce,
        name: "multimem.ld_reduce",
        syntax: "multimem.ld_reduce{.ldsem}{.scope}{.ss}.op.type d, [a];",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::MultimemLdReduce,
        name: "multimem.ld_reduce",
        syntax: "multimem.ld_reduce{.ldsem}{.scope}{.ss}.op{.acc_prec}{.vec}.type d, [a];",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::MultimemRed,
        name: "multimem.red",
        syntax: "multimem.red{.redsem}{.scope}{.ss}.op.type [a], b;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::MultimemRed,
        name: "multimem.red",
        syntax: "multimem.red{.redsem}{.scope}{.ss}.redop{.vec}.redtype [a], b;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::MultimemSt,
        name: "multimem.st",
        syntax: "multimem.st.weak{.ss}.type [a], b;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".weak"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::MultimemSt,
        name: "multimem.st",
        syntax: "multimem.st.weak{.ss}{.vec}.type [a], b;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".weak"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::MultimemSt,
        name: "multimem.st",
        syntax: "multimem.st{.stsem}{.scope}{.ss}.type [a], b;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::MultimemSt,
        name: "multimem.st",
        syntax: "multimem.st{.stsem}{.scope}{.ss}{.vec}.type [a], b;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Nanosleep,
        name: "nanosleep",
        syntax: "nanosleep.u32 t;",
        min_operand_arity: 1usize,
        max_operand_arity: 1usize,
        required_modifiers: &[".u32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Neg,
        name: "neg",
        syntax: "neg.bf16 d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".bf16"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Neg,
        name: "neg",
        syntax: "neg.bf16x2 d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".bf16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Neg,
        name: "neg",
        syntax: "neg.f64 d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".f64"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Neg,
        name: "neg",
        syntax: "neg.type d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Neg,
        name: "neg",
        syntax: "neg{.ftz}.f16 d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".f16"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Neg,
        name: "neg",
        syntax: "neg{.ftz}.f16x2 d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".f16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Neg,
        name: "neg",
        syntax: "neg{.ftz}.f32 d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Not,
        name: "not",
        syntax: "not.type d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Or,
        name: "or",
        syntax: "or.type d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::PmEvent,
        name: "pmevent",
        syntax: "pmevent a; pmevent.mask a;",
        min_operand_arity: 1usize,
        max_operand_arity: 1usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Popc,
        name: "popc",
        syntax: "popc.type d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Prefetch,
        name: "prefetch",
        syntax: "prefetch{.space}.level [a]; prefetch.global.level::eviction_priority [a];",
        min_operand_arity: 1usize,
        max_operand_arity: 1usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Prefetch,
        name: "prefetch",
        syntax: "prefetch{.tensormap_space}.tensormap [a];",
        min_operand_arity: 1usize,
        max_operand_arity: 1usize,
        required_modifiers: &[".tensormap"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Prefetchu,
        name: "prefetchu",
        syntax: "prefetchu.L1 [a];",
        min_operand_arity: 1usize,
        max_operand_arity: 1usize,
        required_modifiers: &[".L1"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Prmt,
        name: "prmt",
        syntax: "prmt.b32{.mode} d, a, b, c;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".b32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Rcp,
        name: "rcp",
        syntax: "rcp.approx{.ftz}.f32 d, a; rcp.rnd{.ftz}.f32 d, a; rcp.rnd.f64 d, a;",
        min_operand_arity: 3usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".approx", ".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::RcpApproxFtzF64,
        name: "rcp.approx.ftz.f64",
        syntax: "rcp.approx.ftz.f64 d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Red,
        name: "red",
        syntax: "red{.sem}{.scope}{.space}.add.noftz{.level::cache_hint}.bf16 [a], b {, cache-policy};",
        min_operand_arity: 2usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".add", ".noftz", ".bf16"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Red,
        name: "red",
        syntax: "red{.sem}{.scope}{.space}.add.noftz{.level::cache_hint}.bf16x2 [a], b {, cache-policy};",
        min_operand_arity: 2usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".add", ".noftz", ".bf16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Red,
        name: "red",
        syntax: "red{.sem}{.scope}{.space}.add.noftz{.level::cache_hint}.f16 [a], b{, cache-policy};",
        min_operand_arity: 2usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".add", ".noftz", ".f16"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Red,
        name: "red",
        syntax: "red{.sem}{.scope}{.space}.add.noftz{.level::cache_hint}.f16x2 [a], b{, cache-policy};",
        min_operand_arity: 2usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".add", ".noftz", ".f16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Red,
        name: "red",
        syntax: "red{.sem}{.scope}{.space}.op{.level::cache_hint}.type [a], b{, cache-policy};",
        min_operand_arity: 2usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::RedAsync,
        name: "red.async",
        syntax: "red.async.sem.scope{.ss}.completion_mechanism.add.type [a], b, [mbar];",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".add"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::RedAsync,
        name: "red.async",
        syntax: "red.async.sem.scope{.ss}.completion_mechanism.op.type [a], b, [mbar];",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::RedAsync,
        name: "red.async",
        syntax: "red.async{.mmio}.sem.scope{.ss}.add.type [a], b;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".add"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::ReduxSync,
        name: "redux.sync",
        syntax: "redux.sync.op.b32 dst, src, membermask;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".b32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::ReduxSync,
        name: "redux.sync",
        syntax: "redux.sync.op.type dst, src, membermask;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::ReduxSync,
        name: "redux.sync",
        syntax: "redux.sync.op{.abs.}{.NaN}.f32 dst, src, membermask;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Rem,
        name: "rem",
        syntax: "rem.type d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Ret,
        name: "ret",
        syntax: "ret{.uni};",
        min_operand_arity: 0usize,
        max_operand_arity: 0usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Rsqrt,
        name: "rsqrt",
        syntax: "rsqrt.approx.f64 d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".approx", ".f64"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Rsqrt,
        name: "rsqrt",
        syntax: "rsqrt.approx{.ftz}.f32 d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".approx", ".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::RsqrtApproxFtzF64,
        name: "rsqrt.approx.ftz.f64",
        syntax: "rsqrt.approx.ftz.f64 d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Sad,
        name: "sad",
        syntax: "sad.type d, a, b, c;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Selp,
        name: "selp",
        syntax: "selp.type d, a, b, c;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Set,
        name: "set",
        syntax: "set.CmpOp.BoolOp.bf16.stype d, a, b, {!}c;",
        min_operand_arity: 3usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".bf16"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Set,
        name: "set",
        syntax: "set.CmpOp.BoolOp.dtype.bf16 d, a, b, {!}c;",
        min_operand_arity: 3usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".bf16"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Set,
        name: "set",
        syntax: "set.CmpOp.BoolOp.dtype.bf16x2 d, a, b, {!}c;",
        min_operand_arity: 3usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".bf16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Set,
        name: "set",
        syntax: "set.CmpOp.BoolOp{.ftz}.dtype.f16 d, a, b, {!}c;",
        min_operand_arity: 3usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".f16"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Set,
        name: "set",
        syntax: "set.CmpOp.BoolOp{.ftz}.dtype.f16x2 d, a, b, {!}c;",
        min_operand_arity: 3usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".f16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Set,
        name: "set",
        syntax: "set.CmpOp.BoolOp{.ftz}.dtype.stype d, a, b, {!}c;",
        min_operand_arity: 3usize,
        max_operand_arity: 4usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Set,
        name: "set",
        syntax: "set.CmpOp.BoolOp{.ftz}.f16.stype d, a, b, {!}c;",
        min_operand_arity: 3usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".f16"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Set,
        name: "set",
        syntax: "set.CmpOp.bf16.stype d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".bf16"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Set,
        name: "set",
        syntax: "set.CmpOp.dtype.bf16 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".bf16"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Set,
        name: "set",
        syntax: "set.CmpOp.dtype.bf16x2 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".bf16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Set,
        name: "set",
        syntax: "set.CmpOp{.ftz}.dtype.f16 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".f16"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Set,
        name: "set",
        syntax: "set.CmpOp{.ftz}.dtype.f16x2 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".f16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Set,
        name: "set",
        syntax: "set.CmpOp{.ftz}.dtype.stype d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Set,
        name: "set",
        syntax: "set.CmpOp{.ftz}.f16.stype d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".f16"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::SetMaxNReg,
        name: "setmaxnreg",
        syntax: "setmaxnreg.action.sync.aligned.u32 imm-reg-count;",
        min_operand_arity: 1usize,
        max_operand_arity: 1usize,
        required_modifiers: &[".sync", ".aligned", ".u32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Setp,
        name: "setp",
        syntax: "setp.CmpOp.BoolOp.bf16 p, a, b, {!}c;",
        min_operand_arity: 3usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".bf16"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Setp,
        name: "setp",
        syntax: "setp.CmpOp.BoolOp.bf16x2 p|q, a, b, {!}c;",
        min_operand_arity: 3usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".bf16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Setp,
        name: "setp",
        syntax: "setp.CmpOp.BoolOp{.ftz}.f16 p, a, b, {!}c;",
        min_operand_arity: 3usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".f16"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Setp,
        name: "setp",
        syntax: "setp.CmpOp.BoolOp{.ftz}.f16x2 p|q, a, b, {!}c;",
        min_operand_arity: 3usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".f16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Setp,
        name: "setp",
        syntax: "setp.CmpOp.BoolOp{.ftz}.type p[|q], a, b, {!}c;",
        min_operand_arity: 3usize,
        max_operand_arity: 4usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Setp,
        name: "setp",
        syntax: "setp.CmpOp.bf16 p, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".bf16"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Setp,
        name: "setp",
        syntax: "setp.CmpOp.bf16x2 p|q, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".bf16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Setp,
        name: "setp",
        syntax: "setp.CmpOp{.ftz}.f16 p, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".f16"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Setp,
        name: "setp",
        syntax: "setp.CmpOp{.ftz}.f16x2 p|q, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".f16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Setp,
        name: "setp",
        syntax: "setp.CmpOp{.ftz}.type p[|q], a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Shf,
        name: "shf",
        syntax: "shf.l.mode.b32 d, a, b, c; shf.r.mode.b32 d, a, b, c;",
        min_operand_arity: 7usize,
        max_operand_arity: 7usize,
        required_modifiers: &[".l", ".b32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Shfl,
        name: "shfl",
        syntax: "shfl.mode.b32 d[|p], a, b, c;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".b32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::ShflSync,
        name: "shfl.sync",
        syntax: "shfl.sync.mode.b32 d[|p], a, b, c, membermask;",
        min_operand_arity: 5usize,
        max_operand_arity: 5usize,
        required_modifiers: &[".b32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Shl,
        name: "shl",
        syntax: "shl.type d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Shr,
        name: "shr",
        syntax: "shr.type d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Sin,
        name: "sin",
        syntax: "sin.approx{.ftz}.f32 d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".approx", ".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Slct,
        name: "slct",
        syntax: "slct.dtype.s32 d, a, b, c;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".s32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Slct,
        name: "slct",
        syntax: "slct{.ftz}.dtype.f32 d, a, b, c;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Sqrt,
        name: "sqrt",
        syntax: "sqrt.approx{.ftz}.f32 d, a; sqrt.rnd{.ftz}.f32 d, a; sqrt.rnd.f64 d, a;",
        min_operand_arity: 3usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".approx", ".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::St,
        name: "st",
        syntax: "st.mmio.relaxed.sys{.global}.type [a], b;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".mmio", ".relaxed", ".sys"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::St,
        name: "st",
        syntax: "st.relaxed.scope{.ss}{.level1::eviction_priority}{.level2::eviction_priority}{.level::cache_hint}{.vec}.type [a], b{, cache-policy};",
        min_operand_arity: 2usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".relaxed"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::St,
        name: "st",
        syntax: "st.release.scope{.ss}{.level1::eviction_priority}{.level2::eviction_priority}{.level::cache_hint}{.vec}.type [a], b{, cache-policy};",
        min_operand_arity: 2usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".release"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::St,
        name: "st",
        syntax: "st.volatile{.ss}{.vec}.type [a], b;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".volatile"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::St,
        name: "st",
        syntax: "st{.weak}{.ss}{.cop}{.level::cache_hint}{.vec}.type [a], b{, cache-policy};",
        min_operand_arity: 2usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::St,
        name: "st",
        syntax: "st{.weak}{.ss}{.level1::eviction_priority}{.level2::eviction_priority}{.level::cache_hint}{.vec}.type [a], b{, cache-policy};",
        min_operand_arity: 2usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::StAsync,
        name: "st.async",
        syntax: "st.async{.mmio}.sem.scope{.ss}.type [a], b;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::StAsync,
        name: "st.async",
        syntax: "st.async{.scope}{.ss}.completion_mechanism{.vec}.type [a], b, [mbar];",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::StAsync,
        name: "st.async",
        syntax: "st.async{.weak}{.ss}.completion_mechanism{.vec}.type [a], b, [mbar];",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::StBulk,
        name: "st.bulk",
        syntax: "st.bulk{.weak}{.shared::cta} [a], size, initval;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::StackRestore,
        name: "stackrestore",
        syntax: "stackrestore.type a;",
        min_operand_arity: 1usize,
        max_operand_arity: 1usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::StackSave,
        name: "stacksave",
        syntax: "stacksave.type d;",
        min_operand_arity: 1usize,
        max_operand_arity: 1usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::StMatrix,
        name: "stmatrix",
        syntax: "stmatrix.sync.aligned.shape.num{.trans}{.ss}.type [p], r;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".sync", ".aligned"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Sub,
        name: "sub",
        syntax: "sub.type1 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Sub,
        name: "sub",
        syntax: "sub{.rnd}.bf16 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".bf16"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Sub,
        name: "sub",
        syntax: "sub{.rnd}.bf16x2 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".bf16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Sub,
        name: "sub",
        syntax: "sub{.rnd}.f64 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".f64"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Sub,
        name: "sub",
        syntax: "sub{.rnd}{.ftz}.f32x2 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".f32x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Sub,
        name: "sub",
        syntax: "sub{.rnd}{.ftz}{.sat}.f16 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".f16"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Sub,
        name: "sub",
        syntax: "sub{.rnd}{.ftz}{.sat}.f16x2 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".f16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Sub,
        name: "sub",
        syntax: "sub{.rnd}{.ftz}{.sat}.f32 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Sub,
        name: "sub",
        syntax: "sub{.rnd}{.sat}.f32.atype d, a, c;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Sub,
        name: "sub",
        syntax: "sub{.sat}.type2 d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::SubCc,
        name: "sub.cc",
        syntax: "sub.cc.type d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Subc,
        name: "subc",
        syntax: "subc{.cc}.type d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Suld,
        name: "suld",
        syntax: "suld.b.geom{.cop}.vec.dtype.clamp d, [a, b];",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".b"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Suq,
        name: "suq",
        syntax: "suq.query.b32 d, [a];",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".b32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Sured,
        name: "sured",
        syntax: "sured.b.op.geom.ctype.clamp [a,b],c; sured.p.op.geom.ctype.clamp [a,b],c;",
        min_operand_arity: 5usize,
        max_operand_arity: 5usize,
        required_modifiers: &[".b"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Sust,
        name: "sust",
        syntax: "sust.b.{1d,2d,3d}{.cop}.vec.ctype.clamp [a, b], c; sust.p.{1d,2d,3d}.vec.b32.clamp [a, b], c;",
        min_operand_arity: 6usize,
        max_operand_arity: 7usize,
        required_modifiers: &[".b", "."],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Sust,
        name: "sust",
        syntax: "sust.b.{a1d,a2d}{.cop}.vec.ctype.clamp [a, b], c;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".b", "."],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Szext,
        name: "szext",
        syntax: "szext.mode.type d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Tanh,
        name: "tanh",
        syntax: "tanh.approx.f32 d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".approx", ".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Tanh,
        name: "tanh",
        syntax: "tanh.approx.type d, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".approx"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Tcgen05Alloc,
        name: "tcgen05.alloc",
        syntax: "tcgen05.alloc.cta_group.sync.aligned{.shared::cta}.b32 [dst], nCols;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".sync", ".aligned", ".b32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Tcgen05Commit,
        name: "tcgen05.commit",
        syntax: "tcgen05.commit.cta_group.completion_mechanism{.shared::cluster}{.multicast}.b64 [mbar] {, ctaMask};",
        min_operand_arity: 1usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".b64"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Tcgen05Cp,
        name: "tcgen05.cp",
        syntax: "tcgen05.cp.cta_group.shape{.multicast}{.dst_fmt.src_fmt} [taddr], s-desc;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Tcgen05Dealloc,
        name: "tcgen05.dealloc",
        syntax: "tcgen05.dealloc.cta_group.sync.aligned.b32 taddr, nCols;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".sync", ".aligned", ".b32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Tcgen05Ld,
        name: "tcgen05.ld",
        syntax: "tcgen05.ld.red.sync.aligned.shape3.num.redOp.type r, redval, [taddr];",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".red", ".sync", ".aligned"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Tcgen05Ld,
        name: "tcgen05.ld",
        syntax: "tcgen05.ld.red.sync.aligned.shape3.num.redOp{.abs}{.NaN}.f32 r, redval, [taddr];",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".red", ".sync", ".aligned", ".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Tcgen05Ld,
        name: "tcgen05.ld",
        syntax: "tcgen05.ld.red.sync.aligned.shape4.num.redOp.type r, redval, [taddr], immHalfSplitoff;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".red", ".sync", ".aligned"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Tcgen05Ld,
        name: "tcgen05.ld",
        syntax: "tcgen05.ld.red.sync.aligned.shape4.num.redOp{.abs}{.NaN}.f32 r, redval, [taddr], immHalfSplitoff;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".red", ".sync", ".aligned", ".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Tcgen05Ld,
        name: "tcgen05.ld",
        syntax: "tcgen05.ld.sync.aligned.shape1.num{.pack}.b32 r, [taddr];",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".sync", ".aligned", ".b32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Tcgen05Ld,
        name: "tcgen05.ld",
        syntax: "tcgen05.ld.sync.aligned.shape2.num{.pack}.b32 r, [taddr], immHalfSplitoff;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".sync", ".aligned", ".b32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Tcgen05RelinquishAllocPermit,
        name: "tcgen05.relinquish_alloc_permit",
        syntax: "tcgen05.relinquish_alloc_permit.cta_group.sync.aligned;",
        min_operand_arity: 0usize,
        max_operand_arity: 0usize,
        required_modifiers: &[".sync", ".aligned"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Tcgen05Shift,
        name: "tcgen05.shift",
        syntax: "tcgen05.shift.cta_group.down [taddr];",
        min_operand_arity: 1usize,
        max_operand_arity: 1usize,
        required_modifiers: &[".down"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Tcgen05St,
        name: "tcgen05.st",
        syntax: "tcgen05.st.sync.aligned.shape1.num{.unpack}.b32 [taddr], r;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".sync", ".aligned", ".b32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Tcgen05St,
        name: "tcgen05.st",
        syntax: "tcgen05.st.sync.aligned.shape2.num{.unpack}.b32 [taddr], immHalfSplitoff, r;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".sync", ".aligned", ".b32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::TensorMapCpFenceProxy,
        name: "tensormap.cp_fenceproxy",
        syntax: "tensormap.cp_fenceproxy.cp_qualifiers.fence_qualifiers.sync.aligned [dst], [src], size;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".sync", ".aligned"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::TensormapReplace,
        name: "tensormap.replace",
        syntax: "tensormap.replace.mode.field1{.ss}.b1024.type [addr], new_val;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".b1024"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::TensormapReplace,
        name: "tensormap.replace",
        syntax: "tensormap.replace.mode.field2{.ss}.b1024.type [addr], ord, new_val;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".b1024"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::TensormapReplace,
        name: "tensormap.replace",
        syntax: "tensormap.replace.mode.field3{.ss}.b1024.type [addr], new_val;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".b1024"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Testp,
        name: "testp",
        syntax: "testp.op.type p, a;",
        min_operand_arity: 2usize,
        max_operand_arity: 2usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Tex,
        name: "tex",
        syntax: "tex.base.geom.v2.f16x2.ctype d[|p], [a, {b,} c] {, e} {, f};",
        min_operand_arity: 3usize,
        max_operand_arity: 6usize,
        required_modifiers: &[".base", ".v2", ".f16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Tex,
        name: "tex",
        syntax: "tex.base.geom.v4.dtype.ctype d[|p], [a, {b,} c] {, e} {, f};",
        min_operand_arity: 3usize,
        max_operand_arity: 6usize,
        required_modifiers: &[".base", ".v4"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Tex,
        name: "tex",
        syntax: "tex.geom.v2.f16x2.ctype d[|p], [a, b, c] {, e} {, f};",
        min_operand_arity: 4usize,
        max_operand_arity: 6usize,
        required_modifiers: &[".v2", ".f16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Tex,
        name: "tex",
        syntax: "tex.geom.v2.f16x2.ctype d[|p], [a, c] {, e} {, f};",
        min_operand_arity: 3usize,
        max_operand_arity: 5usize,
        required_modifiers: &[".v2", ".f16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Tex,
        name: "tex",
        syntax: "tex.geom.v4.dtype.ctype d, [a, c] {, e} {, f};",
        min_operand_arity: 3usize,
        max_operand_arity: 5usize,
        required_modifiers: &[".v4"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Tex,
        name: "tex",
        syntax: "tex.geom.v4.dtype.ctype d[|p], [a, b, c] {, e} {, f};",
        min_operand_arity: 4usize,
        max_operand_arity: 6usize,
        required_modifiers: &[".v4"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Tex,
        name: "tex",
        syntax: "tex.grad.geom.v2.f16x2.ctype d[|p], [a, {b,} c], dPdx, dPdy {, e} {, f};",
        min_operand_arity: 5usize,
        max_operand_arity: 8usize,
        required_modifiers: &[".grad", ".v2", ".f16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Tex,
        name: "tex",
        syntax: "tex.grad.geom.v4.dtype.ctype d[|p], [a, {b,} c], dPdx, dPdy {, e} {, f};",
        min_operand_arity: 5usize,
        max_operand_arity: 8usize,
        required_modifiers: &[".grad", ".v4"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Tex,
        name: "tex",
        syntax: "tex.level.geom.v2.f16x2.ctype d[|p], [a, {b,} c], lod {, e} {, f};",
        min_operand_arity: 4usize,
        max_operand_arity: 7usize,
        required_modifiers: &[".level", ".v2", ".f16x2"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Tex,
        name: "tex",
        syntax: "tex.level.geom.v4.dtype.ctype d[|p], [a, {b,} c], lod {, e} {, f};",
        min_operand_arity: 4usize,
        max_operand_arity: 7usize,
        required_modifiers: &[".level", ".v4"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Tld4,
        name: "tld4",
        syntax: "tld4.comp.2d.v4.dtype.f32 d[|p], [a, c] {, e} {, f};",
        min_operand_arity: 3usize,
        max_operand_arity: 5usize,
        required_modifiers: &[".2d", ".v4", ".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Tld4,
        name: "tld4",
        syntax: "tld4.comp.geom.v4.dtype.f32 d[|p], [a, b, c] {, e} {, f};",
        min_operand_arity: 4usize,
        max_operand_arity: 6usize,
        required_modifiers: &[".v4", ".f32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Trap,
        name: "trap",
        syntax: "trap;",
        min_operand_arity: 0usize,
        max_operand_arity: 0usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Txq,
        name: "txq",
        syntax: "txq.tquery.b32 d, [a]; txq.level.tlquery.b32 d, [a], lod; txq.squery.b32 d, [a];",
        min_operand_arity: 5usize,
        max_operand_arity: 5usize,
        required_modifiers: &[".b32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::VMad,
        name: "vmad",
        syntax: "vmad.dtype.atype.btype.po{.sat}{.scale} d, a{.asel}, b{.bsel}, c;",
        min_operand_arity: 2usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".po"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::VMad,
        name: "vmad",
        syntax: "vmad.dtype.atype.btype{.sat}{.scale} d, {-}a{.asel}, {-}b{.bsel}, {-}c;",
        min_operand_arity: 1usize,
        max_operand_arity: 4usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Vote,
        name: "vote",
        syntax: "vote.ballot.b32 d, {!}a;",
        min_operand_arity: 1usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".ballot", ".b32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Vote,
        name: "vote",
        syntax: "vote.mode.pred d, {!}a;",
        min_operand_arity: 1usize,
        max_operand_arity: 2usize,
        required_modifiers: &[".pred"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::VoteSync,
        name: "vote.sync",
        syntax: "vote.sync.ballot.b32 d, {!}a, membermask;",
        min_operand_arity: 2usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".ballot", ".b32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::VoteSync,
        name: "vote.sync",
        syntax: "vote.sync.mode.pred d, {!}a, membermask;",
        min_operand_arity: 2usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".pred"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::VSet,
        name: "vset",
        syntax: "vset.atype.btype.cmp d, a{.asel}, b{.bsel};",
        min_operand_arity: 1usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::VSet,
        name: "vset",
        syntax: "vset.atype.btype.cmp d.dsel, a{.asel}, b{.bsel}, c;",
        min_operand_arity: 2usize,
        max_operand_arity: 4usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::VSet,
        name: "vset",
        syntax: "vset.atype.btype.cmp.op2 d, a{.asel}, b{.bsel}, c;",
        min_operand_arity: 2usize,
        max_operand_arity: 4usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::VSet2,
        name: "vset2",
        syntax: "vset2.atype.btype.cmp d{.mask}, a{.asel}, b{.bsel}, c;",
        min_operand_arity: 1usize,
        max_operand_arity: 4usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::VSet2,
        name: "vset2",
        syntax: "vset2.atype.btype.cmp.add d{.mask}, a{.asel}, b{.bsel}, c;",
        min_operand_arity: 1usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".add"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::VSet4,
        name: "vset4",
        syntax: "vset4.atype.btype.cmp d{.mask}, a{.asel}, b{.bsel}, c;",
        min_operand_arity: 1usize,
        max_operand_arity: 4usize,
        required_modifiers: &[],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::VSet4,
        name: "vset4",
        syntax: "vset4.atype.btype.cmp.add d{.mask}, a{.asel}, b{.bsel}, c;",
        min_operand_arity: 1usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".add"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::WgmmaCommitGroup,
        name: "wgmma.commit_group",
        syntax: "wgmma.commit_group.sync.aligned;",
        min_operand_arity: 0usize,
        max_operand_arity: 0usize,
        required_modifiers: &[".sync", ".aligned"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::WgmmaFence,
        name: "wgmma.fence",
        syntax: "wgmma.fence.sync.aligned;",
        min_operand_arity: 0usize,
        max_operand_arity: 0usize,
        required_modifiers: &[".sync", ".aligned"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::WgmmaMmaAsync,
        name: "wgmma.mma_async",
        syntax: "wgmma.mma_async.sync.aligned.shape.dtype.f16.f16 d, a, b-desc, scale-d, imm-scale-a, imm-scale-b, imm-trans-b;",
        min_operand_arity: 7usize,
        max_operand_arity: 7usize,
        required_modifiers: &[".sync", ".aligned", ".f16", ".f16"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::WgmmaMmaAsync,
        name: "wgmma.mma_async",
        syntax: "wgmma.mma_async.sync.aligned.shape.dtype.f16.f16 d, a-desc, b-desc, scale-d, imm-scale-a, imm-scale-b, imm-trans-a, imm-trans-b;",
        min_operand_arity: 8usize,
        max_operand_arity: 8usize,
        required_modifiers: &[".sync", ".aligned", ".f16", ".f16"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::WgmmaMmaAsyncSp,
        name: "wgmma.mma_async.sp",
        syntax: "wgmma.mma_async.sp.sync.aligned.shape.dtype.f16.f16 d, a, b-desc, sp-meta, sp-sel, scale-d, imm-scale-a, imm-scale-b, imm-trans-b;",
        min_operand_arity: 9usize,
        max_operand_arity: 9usize,
        required_modifiers: &[".sync", ".aligned", ".f16", ".f16"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::WgmmaMmaAsyncSp,
        name: "wgmma.mma_async.sp",
        syntax: "wgmma.mma_async.sp.sync.aligned.shape.dtype.f16.f16 d, a-desc, b-desc, sp-meta, sp-sel, scale-d, imm-scale-a, imm-scale-b, imm-trans-a, imm-trans-b;",
        min_operand_arity: 10usize,
        max_operand_arity: 10usize,
        required_modifiers: &[".sync", ".aligned", ".f16", ".f16"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::WgmmaWaitGroup,
        name: "wgmma.wait_group",
        syntax: "wgmma.wait_group.sync.aligned N;",
        min_operand_arity: 1usize,
        max_operand_arity: 1usize,
        required_modifiers: &[".sync", ".aligned"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::WmmaLoad,
        name: "wmma.load",
        syntax: "wmma.load.a.sync.aligned.layout.shape{.ss}.atype r, [p] {, stride};",
        min_operand_arity: 2usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".a", ".sync", ".aligned"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::WmmaLoad,
        name: "wmma.load",
        syntax: "wmma.load.b.sync.aligned.layout.shape{.ss}.btype r, [p] {, stride};",
        min_operand_arity: 2usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".b", ".sync", ".aligned"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::WmmaLoad,
        name: "wmma.load",
        syntax: "wmma.load.c.sync.aligned.layout.shape{.ss}.ctype r, [p] {, stride};",
        min_operand_arity: 2usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".c", ".sync", ".aligned"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::WmmaMma,
        name: "wmma.mma",
        syntax: "wmma.mma.sync.aligned.alayout.blayout.shape.dtype.ctype d, a, b, c;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".sync", ".aligned"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::WmmaMma,
        name: "wmma.mma",
        syntax: "wmma.mma.sync.aligned.alayout.blayout.shape.s32.atype.btype.s32{.satfinite} d, a, b, c;",
        min_operand_arity: 4usize,
        max_operand_arity: 4usize,
        required_modifiers: &[".sync", ".aligned", ".s32", ".s32"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::WmmaStore,
        name: "wmma.store",
        syntax: "wmma.store.d.sync.aligned.layout.shape{.ss}.type [p], r {, stride}",
        min_operand_arity: 2usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".d", ".sync", ".aligned"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::WmmaStore,
        name: "wmma.store",
        syntax: "wmma.store.d.sync.aligned.layout.shape{.ss}.type [p], r {, stride};",
        min_operand_arity: 2usize,
        max_operand_arity: 3usize,
        required_modifiers: &[".d", ".sync", ".aligned"],
    },
    InstructionSyntax {
        opcode: InstructionOpcode::Xor,
        name: "xor",
        syntax: "xor.type d, a, b;",
        min_operand_arity: 3usize,
        max_operand_arity: 3usize,
        required_modifiers: &[],
    },
];
pub fn instruction_opcode(name: &str) -> Option<InstructionOpcode> {
    match name {
        "abs" => Some(InstructionOpcode::Abs),
        "activemask" => Some(InstructionOpcode::ActiveMask),
        "add" => Some(InstructionOpcode::Add),
        "add.cc" => Some(InstructionOpcode::AddCc),
        "addc" => Some(InstructionOpcode::Addc),
        "alloca" => Some(InstructionOpcode::Alloca),
        "and" => Some(InstructionOpcode::And),
        "applypriority" => Some(InstructionOpcode::ApplyPriority),
        "atom" => Some(InstructionOpcode::Atom),
        "bar" => Some(InstructionOpcode::Bar),
        "bar.warp.sync" => Some(InstructionOpcode::BarWarpSync),
        "barrier" => Some(InstructionOpcode::Barrier),
        "barrier.cluster" => Some(InstructionOpcode::BarrierCluster),
        "bfe" => Some(InstructionOpcode::Bfe),
        "bfi" => Some(InstructionOpcode::Bfi),
        "bfind" => Some(InstructionOpcode::Bfind),
        "bmsk" => Some(InstructionOpcode::Bmsk),
        "bra" => Some(InstructionOpcode::Bra),
        "brev" => Some(InstructionOpcode::Brev),
        "brkpt" => Some(InstructionOpcode::Brkpt),
        "brx.idx" => Some(InstructionOpcode::BrxIdx),
        "call" => Some(InstructionOpcode::Call),
        "clusterlaunchcontrol.query_cancel" => {
            Some(InstructionOpcode::ClusterLaunchControlQueryCancel)
        }
        "clusterlaunchcontrol.try_cancel" => Some(InstructionOpcode::ClusterLaunchControlTryCancel),
        "clz" => Some(InstructionOpcode::Clz),
        "cnot" => Some(InstructionOpcode::Cnot),
        "copysign" => Some(InstructionOpcode::Copysign),
        "cos" => Some(InstructionOpcode::Cos),
        "cp.async.mbarrier.arrive" => Some(InstructionOpcode::CpAsyncMbarrierArrive),
        "createpolicy" => Some(InstructionOpcode::CreatePolicy),
        "cvt" => Some(InstructionOpcode::Cvt),
        "cvt.pack" => Some(InstructionOpcode::CvtPack),
        "cvta" => Some(InstructionOpcode::Cvta),
        "discard" => Some(InstructionOpcode::Discard),
        "div" => Some(InstructionOpcode::Div),
        "dp2a" => Some(InstructionOpcode::Dp2A),
        "dp4a" => Some(InstructionOpcode::Dp4A),
        "elect.sync" => Some(InstructionOpcode::ElectSync),
        "ex2" => Some(InstructionOpcode::Ex2),
        "exit" => Some(InstructionOpcode::Exit),
        "fence" => Some(InstructionOpcode::Fence),
        "fma" => Some(InstructionOpcode::Fma),
        "fns" => Some(InstructionOpcode::Fns),
        "getctarank" => Some(InstructionOpcode::GetCtaRank),
        "griddepcontrol" => Some(InstructionOpcode::GridDepControl),
        "isspacep" => Some(InstructionOpcode::IsSpaceP),
        "istypep" => Some(InstructionOpcode::IsTypeP),
        "ld" => Some(InstructionOpcode::Ld),
        "ld.global.nc" => Some(InstructionOpcode::LdGlobalNc),
        "ldmatrix" => Some(InstructionOpcode::LdMatrix),
        "ldu" => Some(InstructionOpcode::Ldu),
        "lg2" => Some(InstructionOpcode::Lg2),
        "lop3" => Some(InstructionOpcode::Lop3),
        "mad" => Some(InstructionOpcode::Mad),
        "mad.cc" => Some(InstructionOpcode::MadCc),
        "mad24" => Some(InstructionOpcode::Mad24),
        "madc" => Some(InstructionOpcode::Madc),
        "mapa" => Some(InstructionOpcode::Mapa),
        "match.sync" => Some(InstructionOpcode::MatchSync),
        "max" => Some(InstructionOpcode::Max),
        "mbarrier.arrive" => Some(InstructionOpcode::MbarrierArrive),
        "mbarrier.arrive_drop" => Some(InstructionOpcode::MbarrierArriveDrop),
        "mbarrier.complete_tx" => Some(InstructionOpcode::MbarrierCompleteTx),
        "mbarrier.expect_tx" => Some(InstructionOpcode::MbarrierExpectTx),
        "mbarrier.init" => Some(InstructionOpcode::MbarrierInit),
        "mbarrier.inval" => Some(InstructionOpcode::MbarrierInval),
        "mbarrier.pending_count" => Some(InstructionOpcode::MbarrierPendingCount),
        "mbarrier.test_wait" => Some(InstructionOpcode::MbarrierTestWait),
        "mbarrier.try_wait" => Some(InstructionOpcode::MbarrierTryWait),
        "membar" => Some(InstructionOpcode::Membar),
        "min" => Some(InstructionOpcode::Min),
        "mma" => Some(InstructionOpcode::Mma),
        "mma.sp" => Some(InstructionOpcode::MmaSp),
        "mma.sp::ordered_metadata" => Some(InstructionOpcode::MmaSpOrderedMetadata),
        "mov" => Some(InstructionOpcode::Mov),
        "movmatrix" => Some(InstructionOpcode::MovMatrix),
        "mul" => Some(InstructionOpcode::Mul),
        "mul24" => Some(InstructionOpcode::Mul24),
        "multimem.ld_reduce" => Some(InstructionOpcode::MultimemLdReduce),
        "multimem.red" => Some(InstructionOpcode::MultimemRed),
        "multimem.st" => Some(InstructionOpcode::MultimemSt),
        "nanosleep" => Some(InstructionOpcode::Nanosleep),
        "neg" => Some(InstructionOpcode::Neg),
        "not" => Some(InstructionOpcode::Not),
        "or" => Some(InstructionOpcode::Or),
        "pmevent" => Some(InstructionOpcode::PmEvent),
        "popc" => Some(InstructionOpcode::Popc),
        "prefetch" => Some(InstructionOpcode::Prefetch),
        "prefetchu" => Some(InstructionOpcode::Prefetchu),
        "prmt" => Some(InstructionOpcode::Prmt),
        "rcp" => Some(InstructionOpcode::Rcp),
        "rcp.approx.ftz.f64" => Some(InstructionOpcode::RcpApproxFtzF64),
        "red" => Some(InstructionOpcode::Red),
        "red.async" => Some(InstructionOpcode::RedAsync),
        "redux.sync" => Some(InstructionOpcode::ReduxSync),
        "rem" => Some(InstructionOpcode::Rem),
        "ret" => Some(InstructionOpcode::Ret),
        "rsqrt" => Some(InstructionOpcode::Rsqrt),
        "rsqrt.approx.ftz.f64" => Some(InstructionOpcode::RsqrtApproxFtzF64),
        "sad" => Some(InstructionOpcode::Sad),
        "selp" => Some(InstructionOpcode::Selp),
        "set" => Some(InstructionOpcode::Set),
        "setmaxnreg" => Some(InstructionOpcode::SetMaxNReg),
        "setp" => Some(InstructionOpcode::Setp),
        "shf" => Some(InstructionOpcode::Shf),
        "shfl" => Some(InstructionOpcode::Shfl),
        "shfl.sync" => Some(InstructionOpcode::ShflSync),
        "shl" => Some(InstructionOpcode::Shl),
        "shr" => Some(InstructionOpcode::Shr),
        "sin" => Some(InstructionOpcode::Sin),
        "slct" => Some(InstructionOpcode::Slct),
        "sqrt" => Some(InstructionOpcode::Sqrt),
        "st" => Some(InstructionOpcode::St),
        "st.async" => Some(InstructionOpcode::StAsync),
        "st.bulk" => Some(InstructionOpcode::StBulk),
        "stackrestore" => Some(InstructionOpcode::StackRestore),
        "stacksave" => Some(InstructionOpcode::StackSave),
        "stmatrix" => Some(InstructionOpcode::StMatrix),
        "sub" => Some(InstructionOpcode::Sub),
        "sub.cc" => Some(InstructionOpcode::SubCc),
        "subc" => Some(InstructionOpcode::Subc),
        "suld" => Some(InstructionOpcode::Suld),
        "suq" => Some(InstructionOpcode::Suq),
        "sured" => Some(InstructionOpcode::Sured),
        "sust" => Some(InstructionOpcode::Sust),
        "szext" => Some(InstructionOpcode::Szext),
        "tanh" => Some(InstructionOpcode::Tanh),
        "tcgen05.alloc" => Some(InstructionOpcode::Tcgen05Alloc),
        "tcgen05.commit" => Some(InstructionOpcode::Tcgen05Commit),
        "tcgen05.cp" => Some(InstructionOpcode::Tcgen05Cp),
        "tcgen05.dealloc" => Some(InstructionOpcode::Tcgen05Dealloc),
        "tcgen05.fence" => Some(InstructionOpcode::Tcgen05Fence),
        "tcgen05.ld" => Some(InstructionOpcode::Tcgen05Ld),
        "tcgen05.relinquish_alloc_permit" => Some(InstructionOpcode::Tcgen05RelinquishAllocPermit),
        "tcgen05.shift" => Some(InstructionOpcode::Tcgen05Shift),
        "tcgen05.st" => Some(InstructionOpcode::Tcgen05St),
        "tcgen05.wait" => Some(InstructionOpcode::Tcgen05Wait),
        "tensormap.cp_fenceproxy" => Some(InstructionOpcode::TensorMapCpFenceProxy),
        "tensormap.replace" => Some(InstructionOpcode::TensormapReplace),
        "testp" => Some(InstructionOpcode::Testp),
        "tex" => Some(InstructionOpcode::Tex),
        "tld4" => Some(InstructionOpcode::Tld4),
        "trap" => Some(InstructionOpcode::Trap),
        "txq" => Some(InstructionOpcode::Txq),
        "vabsdiff" => Some(InstructionOpcode::VAbsDiff),
        "vabsdiff2" => Some(InstructionOpcode::VAbsDiff2),
        "vabsdiff4" => Some(InstructionOpcode::VAbsDiff4),
        "vadd" => Some(InstructionOpcode::VAdd),
        "vadd2" => Some(InstructionOpcode::VAdd2),
        "vadd4" => Some(InstructionOpcode::VAdd4),
        "vavrg2" => Some(InstructionOpcode::VAvrg2),
        "vavrg4" => Some(InstructionOpcode::VAvrg4),
        "vmad" => Some(InstructionOpcode::VMad),
        "vmax" => Some(InstructionOpcode::VMax),
        "vmax2" => Some(InstructionOpcode::VMax2),
        "vmax4" => Some(InstructionOpcode::VMax4),
        "vmin" => Some(InstructionOpcode::VMin),
        "vmin2" => Some(InstructionOpcode::VMin2),
        "vmin4" => Some(InstructionOpcode::VMin4),
        "vote" => Some(InstructionOpcode::Vote),
        "vote.sync" => Some(InstructionOpcode::VoteSync),
        "vset" => Some(InstructionOpcode::VSet),
        "vset2" => Some(InstructionOpcode::VSet2),
        "vset4" => Some(InstructionOpcode::VSet4),
        "vshl" => Some(InstructionOpcode::VShl),
        "vshr" => Some(InstructionOpcode::VShr),
        "vsub" => Some(InstructionOpcode::VSub),
        "vsub2" => Some(InstructionOpcode::VSub2),
        "vsub4" => Some(InstructionOpcode::VSub4),
        "wgmma.commit_group" => Some(InstructionOpcode::WgmmaCommitGroup),
        "wgmma.fence" => Some(InstructionOpcode::WgmmaFence),
        "wgmma.mma_async" => Some(InstructionOpcode::WgmmaMmaAsync),
        "wgmma.mma_async.sp" => Some(InstructionOpcode::WgmmaMmaAsyncSp),
        "wgmma.wait_group" => Some(InstructionOpcode::WgmmaWaitGroup),
        "wmma.load" => Some(InstructionOpcode::WmmaLoad),
        "wmma.mma" => Some(InstructionOpcode::WmmaMma),
        "wmma.store" => Some(InstructionOpcode::WmmaStore),
        "xor" => Some(InstructionOpcode::Xor),
        _ => None,
    }
}
pub fn syntax_for_opcode(name: &str) -> &'static [InstructionSyntax] {
    match name {
        "abs" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::Abs,
                name: "abs",
                syntax: "abs.bf16 d, a;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".bf16"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Abs,
                name: "abs",
                syntax: "abs.bf16x2 d, a;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".bf16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Abs,
                name: "abs",
                syntax: "abs.f64 d, a;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".f64"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Abs,
                name: "abs",
                syntax: "abs.type d, a;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Abs,
                name: "abs",
                syntax: "abs{.ftz}.f16 d, a;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".f16"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Abs,
                name: "abs",
                syntax: "abs{.ftz}.f16x2 d, a;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".f16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Abs,
                name: "abs",
                syntax: "abs{.ftz}.f32 d, a;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".f32"],
            },
        ],
        "activemask" => &[InstructionSyntax {
            opcode: InstructionOpcode::ActiveMask,
            name: "activemask",
            syntax: "activemask.b32 d;",
            min_operand_arity: 1usize,
            max_operand_arity: 1usize,
            required_modifiers: &[".b32"],
        }],
        "add" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::Add,
                name: "add",
                syntax: "add.type1 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Add,
                name: "add",
                syntax: "add{.rnd}.bf16 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".bf16"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Add,
                name: "add",
                syntax: "add{.rnd}.bf16x2 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".bf16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Add,
                name: "add",
                syntax: "add{.rnd}.f64 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".f64"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Add,
                name: "add",
                syntax: "add{.rnd}{.ftz}.f32x2 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".f32x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Add,
                name: "add",
                syntax: "add{.rnd}{.ftz}{.sat}.f16 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".f16"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Add,
                name: "add",
                syntax: "add{.rnd}{.ftz}{.sat}.f16x2 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".f16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Add,
                name: "add",
                syntax: "add{.rnd}{.ftz}{.sat}.f32 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".f32"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Add,
                name: "add",
                syntax: "add{.rnd}{.sat}.f32.atype d, a, c;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".f32"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Add,
                name: "add",
                syntax: "add{.sat}.type2 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[],
            },
        ],
        "add.cc" => &[InstructionSyntax {
            opcode: InstructionOpcode::AddCc,
            name: "add.cc",
            syntax: "add.cc.type d, a, b;",
            min_operand_arity: 3usize,
            max_operand_arity: 3usize,
            required_modifiers: &[],
        }],
        "addc" => &[InstructionSyntax {
            opcode: InstructionOpcode::Addc,
            name: "addc",
            syntax: "addc{.cc}.type d, a, b;",
            min_operand_arity: 3usize,
            max_operand_arity: 3usize,
            required_modifiers: &[],
        }],
        "alloca" => &[InstructionSyntax {
            opcode: InstructionOpcode::Alloca,
            name: "alloca",
            syntax: "alloca.type ptr, size{, immAlign};",
            min_operand_arity: 2usize,
            max_operand_arity: 3usize,
            required_modifiers: &[],
        }],
        "and" => &[InstructionSyntax {
            opcode: InstructionOpcode::And,
            name: "and",
            syntax: "and.type d, a, b;",
            min_operand_arity: 3usize,
            max_operand_arity: 3usize,
            required_modifiers: &[],
        }],
        "applypriority" => &[InstructionSyntax {
            opcode: InstructionOpcode::ApplyPriority,
            name: "applypriority",
            syntax: "applypriority{.global}.level::eviction_priority [a], size;",
            min_operand_arity: 2usize,
            max_operand_arity: 2usize,
            required_modifiers: &[],
        }],
        "atom" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::Atom,
                name: "atom",
                syntax: "atom{.sem}{.scope}{.space}.add.noftz{.level::cache_hint}.bf16 d, [a], b{, cache-policy};",
                min_operand_arity: 3usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".add", ".noftz", ".bf16"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Atom,
                name: "atom",
                syntax: "atom{.sem}{.scope}{.space}.add.noftz{.level::cache_hint}.bf16x2 d, [a], b{, cache-policy};",
                min_operand_arity: 3usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".add", ".noftz", ".bf16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Atom,
                name: "atom",
                syntax: "atom{.sem}{.scope}{.space}.add.noftz{.level::cache_hint}.f16 d, [a], b{, cache-policy};",
                min_operand_arity: 3usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".add", ".noftz", ".f16"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Atom,
                name: "atom",
                syntax: "atom{.sem}{.scope}{.space}.add.noftz{.level::cache_hint}.f16x2 d, [a], b{, cache-policy};",
                min_operand_arity: 3usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".add", ".noftz", ".f16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Atom,
                name: "atom",
                syntax: "atom{.sem}{.scope}{.space}.cas.b128 d, [a], b, c;",
                min_operand_arity: 4usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".cas", ".b128"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Atom,
                name: "atom",
                syntax: "atom{.sem}{.scope}{.space}.cas.b16 d, [a], b, c;",
                min_operand_arity: 4usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".cas", ".b16"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Atom,
                name: "atom",
                syntax: "atom{.sem}{.scope}{.space}.exch{.level::cache_hint}.b128 d, [a], b {, cache-policy};",
                min_operand_arity: 3usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".exch", ".b128"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Atom,
                name: "atom",
                syntax: "atom{.sem}{.scope}{.space}.op.type d, [a], b, c;",
                min_operand_arity: 4usize,
                max_operand_arity: 4usize,
                required_modifiers: &[],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Atom,
                name: "atom",
                syntax: "atom{.sem}{.scope}{.space}.op{.level::cache_hint}.type d, [a], b{, cache-policy};",
                min_operand_arity: 3usize,
                max_operand_arity: 4usize,
                required_modifiers: &[],
            },
        ],
        "bar" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::Bar,
                name: "bar",
                syntax: "bar{.cta}.arrive a, b;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".arrive"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Bar,
                name: "bar",
                syntax: "bar{.cta}.red.op.pred p, a{, b}, {!}c;",
                min_operand_arity: 2usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".red", ".pred"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Bar,
                name: "bar",
                syntax: "bar{.cta}.red.popc.u32 d, a{, b}, {!}c;",
                min_operand_arity: 2usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".red", ".popc", ".u32"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Bar,
                name: "bar",
                syntax: "bar{.cta}.sync a{, b};",
                min_operand_arity: 1usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".sync"],
            },
        ],
        "bar.warp.sync" => &[InstructionSyntax {
            opcode: InstructionOpcode::BarWarpSync,
            name: "bar.warp.sync",
            syntax: "bar.warp.sync membermask;",
            min_operand_arity: 1usize,
            max_operand_arity: 1usize,
            required_modifiers: &[],
        }],
        "barrier" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::Barrier,
                name: "barrier",
                syntax: "barrier{.cta}.arrive{.aligned} a, b;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".arrive"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Barrier,
                name: "barrier",
                syntax: "barrier{.cta}.red.op{.aligned}.pred p, a{, b}, {!}c;",
                min_operand_arity: 2usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".red", ".pred"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Barrier,
                name: "barrier",
                syntax: "barrier{.cta}.red.popc{.aligned}.u32 d, a{, b}, {!}c;",
                min_operand_arity: 2usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".red", ".popc", ".u32"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Barrier,
                name: "barrier",
                syntax: "barrier{.cta}.sync{.aligned} a{, b};",
                min_operand_arity: 1usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".sync"],
            },
        ],
        "barrier.cluster" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::BarrierCluster,
                name: "barrier.cluster",
                syntax: "barrier.cluster.arrive{.sem}{.aligned};",
                min_operand_arity: 0usize,
                max_operand_arity: 0usize,
                required_modifiers: &[".arrive"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::BarrierCluster,
                name: "barrier.cluster",
                syntax: "barrier.cluster.wait{.acquire}{.aligned};",
                min_operand_arity: 0usize,
                max_operand_arity: 0usize,
                required_modifiers: &[".wait"],
            },
        ],
        "bfe" => &[InstructionSyntax {
            opcode: InstructionOpcode::Bfe,
            name: "bfe",
            syntax: "bfe.type d, a, b, c;",
            min_operand_arity: 4usize,
            max_operand_arity: 4usize,
            required_modifiers: &[],
        }],
        "bfi" => &[InstructionSyntax {
            opcode: InstructionOpcode::Bfi,
            name: "bfi",
            syntax: "bfi.type f, a, b, c, d;",
            min_operand_arity: 5usize,
            max_operand_arity: 5usize,
            required_modifiers: &[],
        }],
        "bfind" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::Bfind,
                name: "bfind",
                syntax: "bfind.shiftamt.type d, a;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".shiftamt"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Bfind,
                name: "bfind",
                syntax: "bfind.type d, a;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[],
            },
        ],
        "bmsk" => &[InstructionSyntax {
            opcode: InstructionOpcode::Bmsk,
            name: "bmsk",
            syntax: "bmsk.mode.b32 d, a, b;",
            min_operand_arity: 3usize,
            max_operand_arity: 3usize,
            required_modifiers: &[".b32"],
        }],
        "bra" => &[],
        "brev" => &[InstructionSyntax {
            opcode: InstructionOpcode::Brev,
            name: "brev",
            syntax: "brev.type d, a;",
            min_operand_arity: 2usize,
            max_operand_arity: 2usize,
            required_modifiers: &[],
        }],
        "brkpt" => &[InstructionSyntax {
            opcode: InstructionOpcode::Brkpt,
            name: "brkpt",
            syntax: "brkpt;",
            min_operand_arity: 0usize,
            max_operand_arity: 0usize,
            required_modifiers: &[],
        }],
        "brx.idx" => &[InstructionSyntax {
            opcode: InstructionOpcode::BrxIdx,
            name: "brx.idx",
            syntax: "brx.idx{.uni} index, tlist;",
            min_operand_arity: 2usize,
            max_operand_arity: 2usize,
            required_modifiers: &[],
        }],
        "call" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::Call,
                name: "call",
                syntax: "call{.uni} (ret-param), fptr, (param-list), flist;",
                min_operand_arity: 4usize,
                max_operand_arity: 4usize,
                required_modifiers: &[],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Call,
                name: "call",
                syntax: "call{.uni} (ret-param), fptr, (param-list), fproto;",
                min_operand_arity: 4usize,
                max_operand_arity: 4usize,
                required_modifiers: &[],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Call,
                name: "call",
                syntax: "call{.uni} (ret-param), func, (param-list);",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Call,
                name: "call",
                syntax: "call{.uni} fptr, (param-list), flist;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Call,
                name: "call",
                syntax: "call{.uni} fptr, (param-list), fproto;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Call,
                name: "call",
                syntax: "call{.uni} fptr, flist;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Call,
                name: "call",
                syntax: "call{.uni} fptr, fproto;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Call,
                name: "call",
                syntax: "call{.uni} func, (param-list);",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Call,
                name: "call",
                syntax: "call{.uni} func;",
                min_operand_arity: 1usize,
                max_operand_arity: 1usize,
                required_modifiers: &[],
            },
        ],
        "clusterlaunchcontrol.query_cancel" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::ClusterLaunchControlQueryCancel,
                name: "clusterlaunchcontrol.query_cancel",
                syntax: "clusterlaunchcontrol.query_cancel.get_first_ctaid.v4.b32.b128 {xdim, ydim, zdim, _}, try_cancel_response;",
                min_operand_arity: 4usize,
                max_operand_arity: 5usize,
                required_modifiers: &[".get_first_ctaid", ".v4", ".b32", ".b128"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::ClusterLaunchControlQueryCancel,
                name: "clusterlaunchcontrol.query_cancel",
                syntax: "clusterlaunchcontrol.query_cancel.get_first_ctaid{::dimension}.b32.b128 reg, try_cancel_response;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".get_first_ctaid", ".b32", ".b128"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::ClusterLaunchControlQueryCancel,
                name: "clusterlaunchcontrol.query_cancel",
                syntax: "clusterlaunchcontrol.query_cancel.is_canceled.pred.b128 pred, try_cancel_response;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".is_canceled", ".pred", ".b128"],
            },
        ],
        "clusterlaunchcontrol.try_cancel" => &[InstructionSyntax {
            opcode: InstructionOpcode::ClusterLaunchControlTryCancel,
            name: "clusterlaunchcontrol.try_cancel",
            syntax: "clusterlaunchcontrol.try_cancel.async{.space}.completion_mechanism{.multicast::cluster::all}.b128 [addr], [mbar];",
            min_operand_arity: 2usize,
            max_operand_arity: 2usize,
            required_modifiers: &[".async", ".b128"],
        }],
        "clz" => &[InstructionSyntax {
            opcode: InstructionOpcode::Clz,
            name: "clz",
            syntax: "clz.type d, a;",
            min_operand_arity: 2usize,
            max_operand_arity: 2usize,
            required_modifiers: &[],
        }],
        "cnot" => &[InstructionSyntax {
            opcode: InstructionOpcode::Cnot,
            name: "cnot",
            syntax: "cnot.type d, a;",
            min_operand_arity: 2usize,
            max_operand_arity: 2usize,
            required_modifiers: &[],
        }],
        "copysign" => &[InstructionSyntax {
            opcode: InstructionOpcode::Copysign,
            name: "copysign",
            syntax: "copysign.type d, a, b;",
            min_operand_arity: 3usize,
            max_operand_arity: 3usize,
            required_modifiers: &[],
        }],
        "cos" => &[InstructionSyntax {
            opcode: InstructionOpcode::Cos,
            name: "cos",
            syntax: "cos.approx{.ftz}.f32 d, a;",
            min_operand_arity: 2usize,
            max_operand_arity: 2usize,
            required_modifiers: &[".approx", ".f32"],
        }],
        "cp.async.mbarrier.arrive" => &[InstructionSyntax {
            opcode: InstructionOpcode::CpAsyncMbarrierArrive,
            name: "cp.async.mbarrier.arrive",
            syntax: "cp.async.mbarrier.arrive{.noinc}{.shared{::cta}}.b64 [addr];",
            min_operand_arity: 1usize,
            max_operand_arity: 1usize,
            required_modifiers: &[".b64"],
        }],
        "createpolicy" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::CreatePolicy,
                name: "createpolicy",
                syntax: "createpolicy.cvt.L2.b64 cache-policy, access-property;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".cvt", ".L2", ".b64"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::CreatePolicy,
                name: "createpolicy",
                syntax: "createpolicy.fractional.level::primary_priority{.level::secondary_priority}.b64 cache-policy{, fraction};",
                min_operand_arity: 1usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".fractional", ".b64"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::CreatePolicy,
                name: "createpolicy",
                syntax: "createpolicy.range{.global}.level::primary_priority{.level::secondary_priority}.b64 cache-policy, [a], primary-size, total-size;",
                min_operand_arity: 4usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".range", ".b64"],
            },
        ],
        "cvt" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::Cvt,
                name: "cvt",
                syntax: "cvt.frnd2{.relu}{.satfinite}.bf16.f32 d, a;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".bf16", ".f32"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Cvt,
                name: "cvt",
                syntax: "cvt.frnd2{.relu}{.satfinite}.bf16x2.f32 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".bf16x2", ".f32"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Cvt,
                name: "cvt",
                syntax: "cvt.frnd2{.relu}{.satfinite}.f16.f32 d, a;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".f16", ".f32"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Cvt,
                name: "cvt",
                syntax: "cvt.frnd2{.relu}{.satfinite}.f16x2.f32 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".f16x2", ".f32"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Cvt,
                name: "cvt",
                syntax: "cvt.frnd2{.satfinite}{.relu}.tf32.f32 d, a;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".tf32", ".f32"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Cvt,
                name: "cvt",
                syntax: "cvt.frnd3{.satfinite}.ue8m0x2.bf16x2 d, a;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".ue8m0x2", ".bf16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Cvt,
                name: "cvt",
                syntax: "cvt.frnd3{.satfinite}.ue8m0x2.f32 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".ue8m0x2", ".f32"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Cvt,
                name: "cvt",
                syntax: "cvt.rn.bf16x2.ue8m0x2 d, a;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".rn", ".bf16x2", ".ue8m0x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Cvt,
                name: "cvt",
                syntax: "cvt.rn.satfinite{.relu}.f4x2type.f32 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".rn", ".satfinite", ".f32"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Cvt,
                name: "cvt",
                syntax: "cvt.rn.satfinite{.relu}.f4x2type.fp16x2type d, a;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".rn", ".satfinite"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Cvt,
                name: "cvt",
                syntax: "cvt.rn.satfinite{.relu}.f6x2type.f32 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".rn", ".satfinite", ".f32"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Cvt,
                name: "cvt",
                syntax: "cvt.rn.satfinite{.relu}.f6x2type.fp16x2type d, a;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".rn", ".satfinite"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Cvt,
                name: "cvt",
                syntax: "cvt.rn.satfinite{.relu}.f8x2type.f32 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".rn", ".satfinite", ".f32"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Cvt,
                name: "cvt",
                syntax: "cvt.rn.satfinite{.relu}.f8x2type.fp16x2 d, a;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".rn", ".satfinite", ".fp16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Cvt,
                name: "cvt",
                syntax: "cvt.rn.satfinite{.relu}{.scaled::n2::ue8m0}.s2f6x2.bf16x2 d, a{, scale-factor};",
                min_operand_arity: 2usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".rn", ".satfinite", ".s2f6x2", ".bf16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Cvt,
                name: "cvt",
                syntax: "cvt.rn.satfinite{.relu}{.scaled::n2::ue8m0}.s2f6x2.f32 d, a, b{, scale-factor};",
                min_operand_arity: 3usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".rn", ".satfinite", ".s2f6x2", ".f32"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Cvt,
                name: "cvt",
                syntax: "cvt.rna{.satfinite}.tf32.f32 d, a;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".rna", ".tf32", ".f32"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Cvt,
                name: "cvt",
                syntax: "cvt.rn{.relu}.f16x2.f4x2type d, a;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".rn", ".f16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Cvt,
                name: "cvt",
                syntax: "cvt.rn{.relu}.f16x2.f6x2type d, a;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".rn", ".f16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Cvt,
                name: "cvt",
                syntax: "cvt.rn{.relu}.f16x2.f8x2type d, a;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".rn", ".f16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Cvt,
                name: "cvt",
                syntax: "cvt.rn{.relu}{.satfinite}{.scaled::n2::ue8m0}.bf16x2.f4x2type d, a{, scale-factor};",
                min_operand_arity: 2usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".rn", ".bf16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Cvt,
                name: "cvt",
                syntax: "cvt.rn{.relu}{.satfinite}{.scaled::n2::ue8m0}.bf16x2.f6x2type d, a{, scale-factor};",
                min_operand_arity: 2usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".rn", ".bf16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Cvt,
                name: "cvt",
                syntax: "cvt.rn{.relu}{.satfinite}{.scaled::n2::ue8m0}.bf16x2.f8x2type d, a{, scale-factor};",
                min_operand_arity: 2usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".rn", ".bf16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Cvt,
                name: "cvt",
                syntax: "cvt.rn{.satfinite}{.relu}{.scaled::n2::ue8m0}.bf16x2.s2f6x2 d, a{, scale-factor};",
                min_operand_arity: 2usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".rn", ".bf16x2", ".s2f6x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Cvt,
                name: "cvt",
                syntax: "cvt.rs{.relu}.satfinite.f4x4type.f32 d, {a, b, e, f}, rbits;",
                min_operand_arity: 5usize,
                max_operand_arity: 6usize,
                required_modifiers: &[".rs", ".satfinite", ".f32"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Cvt,
                name: "cvt",
                syntax: "cvt.rs{.relu}.satfinite.f6x4type.f32 d, {a, b, e, f}, rbits;",
                min_operand_arity: 5usize,
                max_operand_arity: 6usize,
                required_modifiers: &[".rs", ".satfinite", ".f32"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Cvt,
                name: "cvt",
                syntax: "cvt.rs{.relu}.satfinite.f8x4type.f32 d, {a, b, e, f}, rbits;",
                min_operand_arity: 5usize,
                max_operand_arity: 6usize,
                required_modifiers: &[".rs", ".satfinite", ".f32"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Cvt,
                name: "cvt",
                syntax: "cvt.rs{.relu}{.satfinite}.bf16x2.f32 d, a, b, rbits;",
                min_operand_arity: 4usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".rs", ".bf16x2", ".f32"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Cvt,
                name: "cvt",
                syntax: "cvt.rs{.relu}{.satfinite}.f16x2.f32 d, a, b, rbits;",
                min_operand_arity: 4usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".rs", ".f16x2", ".f32"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Cvt,
                name: "cvt",
                syntax: "cvt{.irnd}{.ftz}{.sat}.dtype.atype d, a; cvt{.frnd}{.ftz}{.sat}.dtype.atype d, a;",
                min_operand_arity: 2usize,
                max_operand_arity: 3usize,
                required_modifiers: &[],
            },
        ],
        "cvt.pack" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::CvtPack,
                name: "cvt.pack",
                syntax: "cvt.pack.sat.convertType.abType d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".sat"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::CvtPack,
                name: "cvt.pack",
                syntax: "cvt.pack.sat.convertType.abType.cType d, a, b, c;",
                min_operand_arity: 4usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".sat"],
            },
        ],
        "cvta" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::Cvta,
                name: "cvta",
                syntax: "cvta.space.size p, a; cvta.space.size p, var; cvta.space.size p, var+imm;",
                min_operand_arity: 4usize,
                max_operand_arity: 4usize,
                required_modifiers: &[],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Cvta,
                name: "cvta",
                syntax: "cvta.to.space.size p, a;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".to"],
            },
        ],
        "discard" => &[InstructionSyntax {
            opcode: InstructionOpcode::Discard,
            name: "discard",
            syntax: "discard{.global}.level [a], size;",
            min_operand_arity: 2usize,
            max_operand_arity: 2usize,
            required_modifiers: &[],
        }],
        "div" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::Div,
                name: "div",
                syntax: "div.approx{.ftz}.f32 d, a, b; div.full{.ftz}.f32 d, a, b; div.rnd{.ftz}.f32 d, a, b; div.rnd.f64 d, a, b;",
                min_operand_arity: 7usize,
                max_operand_arity: 9usize,
                required_modifiers: &[".approx", ".f32"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Div,
                name: "div",
                syntax: "div.type d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[],
            },
        ],
        "dp2a" => &[InstructionSyntax {
            opcode: InstructionOpcode::Dp2A,
            name: "dp2a",
            syntax: "dp2a.mode.atype.btype d, a, b, c;",
            min_operand_arity: 4usize,
            max_operand_arity: 4usize,
            required_modifiers: &[],
        }],
        "dp4a" => &[InstructionSyntax {
            opcode: InstructionOpcode::Dp4A,
            name: "dp4a",
            syntax: "dp4a.atype.btype d, a, b, c;",
            min_operand_arity: 4usize,
            max_operand_arity: 4usize,
            required_modifiers: &[],
        }],
        "elect.sync" => &[InstructionSyntax {
            opcode: InstructionOpcode::ElectSync,
            name: "elect.sync",
            syntax: "elect.sync d|p, membermask;",
            min_operand_arity: 2usize,
            max_operand_arity: 2usize,
            required_modifiers: &[],
        }],
        "ex2" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::Ex2,
                name: "ex2",
                syntax: "ex2.approx.atype d, a;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".approx"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Ex2,
                name: "ex2",
                syntax: "ex2.approx.ftz.btype d, a;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".approx", ".ftz"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Ex2,
                name: "ex2",
                syntax: "ex2.approx{.ftz}.f32 d, a;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".approx", ".f32"],
            },
        ],
        "exit" => &[InstructionSyntax {
            opcode: InstructionOpcode::Exit,
            name: "exit",
            syntax: "exit;",
            min_operand_arity: 0usize,
            max_operand_arity: 0usize,
            required_modifiers: &[],
        }],
        "fence" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::Fence,
                name: "fence",
                syntax: "fence.acquire.sync_restrict::shared::cluster.cluster;",
                min_operand_arity: 0usize,
                max_operand_arity: 0usize,
                required_modifiers: &[".acquire", ".sync_restrict::shared::cluster", ".cluster"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Fence,
                name: "fence",
                syntax: "fence.op_restrict.release.cluster;",
                min_operand_arity: 0usize,
                max_operand_arity: 0usize,
                required_modifiers: &[".release", ".cluster"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Fence,
                name: "fence",
                syntax: "fence.proxy.async::generic.acquire.sync_restrict::shared::cluster.cluster;",
                min_operand_arity: 0usize,
                max_operand_arity: 0usize,
                required_modifiers: &[
                    ".proxy",
                    ".async::generic",
                    ".acquire",
                    ".sync_restrict::shared::cluster",
                    ".cluster",
                ],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Fence,
                name: "fence",
                syntax: "fence.proxy.async::generic.release.sync_restrict::shared::cta.cluster;",
                min_operand_arity: 0usize,
                max_operand_arity: 0usize,
                required_modifiers: &[
                    ".proxy",
                    ".async::generic",
                    ".release",
                    ".sync_restrict::shared::cta",
                    ".cluster",
                ],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Fence,
                name: "fence",
                syntax: "fence.proxy.proxykind;",
                min_operand_arity: 0usize,
                max_operand_arity: 0usize,
                required_modifiers: &[".proxy"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Fence,
                name: "fence",
                syntax: "fence.proxy.to_proxykind::from_proxykind.acquire.scope [addr], size;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".proxy", ".acquire"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Fence,
                name: "fence",
                syntax: "fence.proxy.to_proxykind::from_proxykind.release.scope;",
                min_operand_arity: 0usize,
                max_operand_arity: 0usize,
                required_modifiers: &[".proxy", ".release"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Fence,
                name: "fence",
                syntax: "fence.release.sync_restrict::shared::cta.cluster;",
                min_operand_arity: 0usize,
                max_operand_arity: 0usize,
                required_modifiers: &[".release", ".sync_restrict::shared::cta", ".cluster"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Fence,
                name: "fence",
                syntax: "fence{.sem}.scope;",
                min_operand_arity: 0usize,
                max_operand_arity: 0usize,
                required_modifiers: &[],
            },
        ],
        "fma" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::Fma,
                name: "fma",
                syntax: "fma.rnd.f64 d, a, b, c;",
                min_operand_arity: 4usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".f64"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Fma,
                name: "fma",
                syntax: "fma.rnd.oob.{relu}.type d, a, b, c;",
                min_operand_arity: 4usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".oob", ".", ".type"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Fma,
                name: "fma",
                syntax: "fma.rnd{.ftz}.f32x2 d, a, b, c;",
                min_operand_arity: 4usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".f32x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Fma,
                name: "fma",
                syntax: "fma.rnd{.ftz}.relu.f16 d, a, b, c;",
                min_operand_arity: 4usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".relu", ".f16"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Fma,
                name: "fma",
                syntax: "fma.rnd{.ftz}.relu.f16x2 d, a, b, c;",
                min_operand_arity: 4usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".relu", ".f16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Fma,
                name: "fma",
                syntax: "fma.rnd{.ftz}{.sat}.f16 d, a, b, c;",
                min_operand_arity: 4usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".f16"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Fma,
                name: "fma",
                syntax: "fma.rnd{.ftz}{.sat}.f16x2 d, a, b, c;",
                min_operand_arity: 4usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".f16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Fma,
                name: "fma",
                syntax: "fma.rnd{.ftz}{.sat}.f32 d, a, b, c;",
                min_operand_arity: 4usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".f32"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Fma,
                name: "fma",
                syntax: "fma.rnd{.relu}.bf16 d, a, b, c;",
                min_operand_arity: 4usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".bf16"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Fma,
                name: "fma",
                syntax: "fma.rnd{.relu}.bf16x2 d, a, b, c;",
                min_operand_arity: 4usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".bf16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Fma,
                name: "fma",
                syntax: "fma.rnd{.sat}.f32.abtype d, a, b, c;",
                min_operand_arity: 4usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".f32"],
            },
        ],
        "fns" => &[InstructionSyntax {
            opcode: InstructionOpcode::Fns,
            name: "fns",
            syntax: "fns.b32 d, mask, base, offset;",
            min_operand_arity: 4usize,
            max_operand_arity: 4usize,
            required_modifiers: &[".b32"],
        }],
        "getctarank" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::GetCtaRank,
                name: "getctarank",
                syntax: "getctarank.shared::cluster.type d, a;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".shared::cluster"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::GetCtaRank,
                name: "getctarank",
                syntax: "getctarank.shared::cluster.type d, var + imm;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".shared::cluster"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::GetCtaRank,
                name: "getctarank",
                syntax: "getctarank.shared::cluster.type d, var;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".shared::cluster"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::GetCtaRank,
                name: "getctarank",
                syntax: "getctarank.type d, a;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::GetCtaRank,
                name: "getctarank",
                syntax: "getctarank{.space}.type d, a;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[],
            },
        ],
        "griddepcontrol" => &[InstructionSyntax {
            opcode: InstructionOpcode::GridDepControl,
            name: "griddepcontrol",
            syntax: "griddepcontrol.action;",
            min_operand_arity: 0usize,
            max_operand_arity: 0usize,
            required_modifiers: &[],
        }],
        "isspacep" => &[InstructionSyntax {
            opcode: InstructionOpcode::IsSpaceP,
            name: "isspacep",
            syntax: "isspacep.space p, a;",
            min_operand_arity: 2usize,
            max_operand_arity: 2usize,
            required_modifiers: &[],
        }],
        "istypep" => &[InstructionSyntax {
            opcode: InstructionOpcode::IsTypeP,
            name: "istypep",
            syntax: "istypep.type p, a;",
            min_operand_arity: 2usize,
            max_operand_arity: 2usize,
            required_modifiers: &[],
        }],
        "ld" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::Ld,
                name: "ld",
                syntax: "ld.acquire.scope{.ss}{.level1::eviction_priority}{.level2::eviction_priority}{.level::cache_hint}{.level::prefetch_size}{.vec}.type d, [a]{, cache-policy};",
                min_operand_arity: 2usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".acquire"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Ld,
                name: "ld",
                syntax: "ld.mmio.relaxed.sys{.global}.type d, [a];",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".mmio", ".relaxed", ".sys"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Ld,
                name: "ld",
                syntax: "ld.relaxed.scope{.ss}{.level1::eviction_priority}{.level2::eviction_priority}{.level::cache_hint}{.level::prefetch_size}{.vec}.type d, [a]{, cache-policy};",
                min_operand_arity: 2usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".relaxed"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Ld,
                name: "ld",
                syntax: "ld.volatile{.ss}{.level::prefetch_size}{.vec}.type d, [a];",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".volatile"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Ld,
                name: "ld",
                syntax: "ld{.weak}{.ss}{.cop}{.level::cache_hint}{.level::prefetch_size}{.vec}.type d, [a]{.unified}{, cache-policy};",
                min_operand_arity: 2usize,
                max_operand_arity: 3usize,
                required_modifiers: &[],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Ld,
                name: "ld",
                syntax: "ld{.weak}{.ss}{.level1::eviction_priority}{.level2::eviction_priority}{.level::cache_hint}{.level::prefetch_size}{.vec}.type d, [a]{.unified}{, cache-policy};",
                min_operand_arity: 2usize,
                max_operand_arity: 3usize,
                required_modifiers: &[],
            },
        ],
        "ld.global.nc" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::LdGlobalNc,
                name: "ld.global.nc",
                syntax: "ld.global.nc{.level1::eviction_priority}{.level2::eviction_priority}{.level::cache_hint}{.level::prefetch_size}.type d, [a]{, cache-policy};",
                min_operand_arity: 2usize,
                max_operand_arity: 3usize,
                required_modifiers: &[],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::LdGlobalNc,
                name: "ld.global.nc",
                syntax: "ld.global.nc{.level1::eviction_priority}{.level2::eviction_priority}{.level::cache_hint}{.level::prefetch_size}.vec.type d, [a]{, cache-policy};",
                min_operand_arity: 2usize,
                max_operand_arity: 3usize,
                required_modifiers: &[],
            },
        ],
        "ldmatrix" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::LdMatrix,
                name: "ldmatrix",
                syntax: "ldmatrix.sync.aligned.m16n16.num.trans{.ss}.dst_fmt.src_fmt r, [p];",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".sync", ".aligned", ".m16n16", ".trans"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::LdMatrix,
                name: "ldmatrix",
                syntax: "ldmatrix.sync.aligned.m8n16.num{.ss}.dst_fmt.src_fmt r, [p];",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".sync", ".aligned", ".m8n16"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::LdMatrix,
                name: "ldmatrix",
                syntax: "ldmatrix.sync.aligned.shape.num{.trans}{.ss}.type r, [p];",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".sync", ".aligned"],
            },
        ],
        "ldu" => &[InstructionSyntax {
            opcode: InstructionOpcode::Ldu,
            name: "ldu",
            syntax: "ldu{.ss}.type d, [a]; ldu{.ss}.vec.type d, [a];",
            min_operand_arity: 2usize,
            max_operand_arity: 3usize,
            required_modifiers: &[],
        }],
        "lg2" => &[InstructionSyntax {
            opcode: InstructionOpcode::Lg2,
            name: "lg2",
            syntax: "lg2.approx{.ftz}.f32 d, a;",
            min_operand_arity: 2usize,
            max_operand_arity: 2usize,
            required_modifiers: &[".approx", ".f32"],
        }],
        "lop3" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::Lop3,
                name: "lop3",
                syntax: "lop3.BoolOp.b32 d|p, a, b, c, immLut, q;",
                min_operand_arity: 6usize,
                max_operand_arity: 6usize,
                required_modifiers: &[".b32"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Lop3,
                name: "lop3",
                syntax: "lop3.b32 d, a, b, c, immLut;",
                min_operand_arity: 5usize,
                max_operand_arity: 5usize,
                required_modifiers: &[".b32"],
            },
        ],
        "mad" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::Mad,
                name: "mad",
                syntax: "mad.hi.sat.s32 d, a, b, c;",
                min_operand_arity: 4usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".hi", ".sat", ".s32"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Mad,
                name: "mad",
                syntax: "mad.mode.type d, a, b, c;",
                min_operand_arity: 4usize,
                max_operand_arity: 4usize,
                required_modifiers: &[],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Mad,
                name: "mad",
                syntax: "mad{.ftz}{.sat}.f32 d, a, b, c; mad.rnd{.ftz}{.sat}.f32 d, a, b, c; mad.rnd.f64 d, a, b, c;",
                min_operand_arity: 9usize,
                max_operand_arity: 10usize,
                required_modifiers: &[".f32"],
            },
        ],
        "mad.cc" => &[],
        "mad24" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::Mad24,
                name: "mad24",
                syntax: "mad24.hi.sat.s32 d, a, b, c;",
                min_operand_arity: 4usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".hi", ".sat", ".s32"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Mad24,
                name: "mad24",
                syntax: "mad24.mode.type d, a, b, c;",
                min_operand_arity: 4usize,
                max_operand_arity: 4usize,
                required_modifiers: &[],
            },
        ],
        "madc" => &[InstructionSyntax {
            opcode: InstructionOpcode::Madc,
            name: "madc",
            syntax: "madc{.hi,.lo}{.cc}.type d, a, b, c;",
            min_operand_arity: 4usize,
            max_operand_arity: 4usize,
            required_modifiers: &[],
        }],
        "mapa" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::Mapa,
                name: "mapa",
                syntax: "mapa.shared::cluster.type d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".shared::cluster"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Mapa,
                name: "mapa",
                syntax: "mapa.shared::cluster.type d, sh + imm, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".shared::cluster"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Mapa,
                name: "mapa",
                syntax: "mapa.shared::cluster.type d, sh, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".shared::cluster"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Mapa,
                name: "mapa",
                syntax: "mapa.type d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Mapa,
                name: "mapa",
                syntax: "mapa{.space}.type d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[],
            },
        ],
        "match.sync" => &[],
        "max" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::Max,
                name: "max",
                syntax: "max.f64 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".f64"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Max,
                name: "max",
                syntax: "max.type1 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Max,
                name: "max",
                syntax: "max{.NaN}{.xorsign.abs}.bf16 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".bf16"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Max,
                name: "max",
                syntax: "max{.NaN}{.xorsign.abs}.bf16x2 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".bf16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Max,
                name: "max",
                syntax: "max{.ftz}{.NaN}{.abs}.f32 d, a, b, c;",
                min_operand_arity: 4usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".f32"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Max,
                name: "max",
                syntax: "max{.ftz}{.NaN}{.xorsign.abs}.f16 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".f16"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Max,
                name: "max",
                syntax: "max{.ftz}{.NaN}{.xorsign.abs}.f16x2 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".f16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Max,
                name: "max",
                syntax: "max{.ftz}{.NaN}{.xorsign.abs}.f32 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".f32"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Max,
                name: "max",
                syntax: "max{.relu}.type2 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[],
            },
        ],
        "mbarrier.arrive" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::MbarrierArrive,
                name: "mbarrier.arrive",
                syntax: "mbarrier.arrive.expect_tx{.sem.scope}{.shared::cluster}.b64 _, [addr], txCount;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".expect_tx", ".b64"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::MbarrierArrive,
                name: "mbarrier.arrive",
                syntax: "mbarrier.arrive.noComplete{.release.cta}{.shared{::cta}}.b64 state, [addr], count;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".noComplete", ".b64"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::MbarrierArrive,
                name: "mbarrier.arrive",
                syntax: "mbarrier.arrive{.sem.scope}{.shared::cluster}.b64 _, [addr] {,count} mbarrier.arrive.expect_tx{.sem.scope}{.shared{::cta}}.b64 state, [addr], txCount;",
                min_operand_arity: 3usize,
                max_operand_arity: 5usize,
                required_modifiers: &[".b64"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::MbarrierArrive,
                name: "mbarrier.arrive",
                syntax: "mbarrier.arrive{.sem.scope}{.shared{::cta}}.b64 state, [addr]{, count};",
                min_operand_arity: 2usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".b64"],
            },
        ],
        "mbarrier.arrive_drop" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::MbarrierArriveDrop,
                name: "mbarrier.arrive_drop",
                syntax: "mbarrier.arrive_drop.expect_tx{.sem.scope}{.shared::cluster}.b64 _, [addr], tx_count;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".expect_tx", ".b64"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::MbarrierArriveDrop,
                name: "mbarrier.arrive_drop",
                syntax: "mbarrier.arrive_drop.expect_tx{.sem.scope}{.shared{::cta}}.b64 state, [addr], tx_count;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".expect_tx", ".b64"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::MbarrierArriveDrop,
                name: "mbarrier.arrive_drop",
                syntax: "mbarrier.arrive_drop.noComplete{.release.cta}{.shared{::cta}}.b64 state, [addr], count;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".noComplete", ".b64"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::MbarrierArriveDrop,
                name: "mbarrier.arrive_drop",
                syntax: "mbarrier.arrive_drop{.sem.scope}{.shared::cluster}.b64 _, [addr] {, count};",
                min_operand_arity: 2usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".b64"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::MbarrierArriveDrop,
                name: "mbarrier.arrive_drop",
                syntax: "mbarrier.arrive_drop{.sem.scope}{.shared{::cta}}.b64 state, [addr] {, count};",
                min_operand_arity: 2usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".b64"],
            },
        ],
        "mbarrier.complete_tx" => &[InstructionSyntax {
            opcode: InstructionOpcode::MbarrierCompleteTx,
            name: "mbarrier.complete_tx",
            syntax: "mbarrier.complete_tx{.sem.scope}{.space}.b64 [addr], txCount;",
            min_operand_arity: 2usize,
            max_operand_arity: 2usize,
            required_modifiers: &[".b64"],
        }],
        "mbarrier.expect_tx" => &[InstructionSyntax {
            opcode: InstructionOpcode::MbarrierExpectTx,
            name: "mbarrier.expect_tx",
            syntax: "mbarrier.expect_tx{.sem.scope}{.space}.b64 [addr], txCount;",
            min_operand_arity: 2usize,
            max_operand_arity: 2usize,
            required_modifiers: &[".b64"],
        }],
        "mbarrier.init" => &[InstructionSyntax {
            opcode: InstructionOpcode::MbarrierInit,
            name: "mbarrier.init",
            syntax: "mbarrier.init{.shared{::cta}}.b64 [addr], count;",
            min_operand_arity: 2usize,
            max_operand_arity: 2usize,
            required_modifiers: &[".b64"],
        }],
        "mbarrier.inval" => &[InstructionSyntax {
            opcode: InstructionOpcode::MbarrierInval,
            name: "mbarrier.inval",
            syntax: "mbarrier.inval{.shared{::cta}}.b64 [addr];",
            min_operand_arity: 1usize,
            max_operand_arity: 1usize,
            required_modifiers: &[".b64"],
        }],
        "mbarrier.pending_count" => &[InstructionSyntax {
            opcode: InstructionOpcode::MbarrierPendingCount,
            name: "mbarrier.pending_count",
            syntax: "mbarrier.pending_count.b64 count, state;",
            min_operand_arity: 2usize,
            max_operand_arity: 2usize,
            required_modifiers: &[".b64"],
        }],
        "mbarrier.test_wait" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::MbarrierTestWait,
                name: "mbarrier.test_wait",
                syntax: "mbarrier.test_wait.parity{.sem.scope}{.shared{::cta}}.b64 waitComplete, [addr], phaseParity;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".parity", ".b64"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::MbarrierTestWait,
                name: "mbarrier.test_wait",
                syntax: "mbarrier.test_wait{.sem.scope}{.shared{::cta}}.b64 waitComplete, [addr], state;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".b64"],
            },
        ],
        "mbarrier.try_wait" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::MbarrierTryWait,
                name: "mbarrier.try_wait",
                syntax: "mbarrier.try_wait.parity{.sem.scope}{.shared{::cta}}.b64 waitComplete, [addr], phaseParity {, timeHint};",
                min_operand_arity: 3usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".parity", ".b64"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::MbarrierTryWait,
                name: "mbarrier.try_wait",
                syntax: "mbarrier.try_wait{.sem.scope}{.shared{::cta}}.b64 waitComplete, [addr], state {, timeHint};",
                min_operand_arity: 3usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".b64"],
            },
        ],
        "membar" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::Membar,
                name: "membar",
                syntax: "membar.level;",
                min_operand_arity: 0usize,
                max_operand_arity: 0usize,
                required_modifiers: &[],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Membar,
                name: "membar",
                syntax: "membar.proxy.proxykind;",
                min_operand_arity: 0usize,
                max_operand_arity: 0usize,
                required_modifiers: &[".proxy"],
            },
        ],
        "min" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::Min,
                name: "min",
                syntax: "min.f64 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".f64"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Min,
                name: "min",
                syntax: "min.type1 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Min,
                name: "min",
                syntax: "min{.NaN}{.xorsign.abs}.bf16 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".bf16"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Min,
                name: "min",
                syntax: "min{.NaN}{.xorsign.abs}.bf16x2 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".bf16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Min,
                name: "min",
                syntax: "min{.ftz}{.NaN}{.abs}.f32 d, a, b, c;",
                min_operand_arity: 4usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".f32"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Min,
                name: "min",
                syntax: "min{.ftz}{.NaN}{.xorsign.abs}.f16 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".f16"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Min,
                name: "min",
                syntax: "min{.ftz}{.NaN}{.xorsign.abs}.f16x2 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".f16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Min,
                name: "min",
                syntax: "min{.ftz}{.NaN}{.xorsign.abs}.f32 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".f32"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Min,
                name: "min",
                syntax: "min{.relu}.type2 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[],
            },
        ],
        "mma" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::Mma,
                name: "mma",
                syntax: "mma.sync.aligned.m16n8k16.row.col.dtype.f16.f16.ctype d, a, b, c;",
                min_operand_arity: 4usize,
                max_operand_arity: 4usize,
                required_modifiers: &[
                    ".sync",
                    ".aligned",
                    ".m16n8k16",
                    ".row",
                    ".col",
                    ".f16",
                    ".f16",
                ],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Mma,
                name: "mma",
                syntax: "mma.sync.aligned.m16n8k8.row.col.dtype.f16.f16.ctype d, a, b, c;",
                min_operand_arity: 4usize,
                max_operand_arity: 4usize,
                required_modifiers: &[
                    ".sync", ".aligned", ".m16n8k8", ".row", ".col", ".f16", ".f16",
                ],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Mma,
                name: "mma",
                syntax: "mma.sync.aligned.m8n8k4.alayout.blayout.dtype.f16.f16.ctype d, a, b, c;",
                min_operand_arity: 4usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".sync", ".aligned", ".m8n8k4", ".f16", ".f16"],
            },
        ],
        "mma.sp" => &[],
        "mma.sp::ordered_metadata" => &[],
        "mov" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::Mov,
                name: "mov",
                syntax: "mov.type d, a;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Mov,
                name: "mov",
                syntax: "mov.type d, avar; mov.type d, avar+imm; mov.u32 d, fname; mov.u64 d, fname; mov.u32 d, kernel; mov.u64 d, kernel;",
                min_operand_arity: 7usize,
                max_operand_arity: 7usize,
                required_modifiers: &[],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Mov,
                name: "mov",
                syntax: "mov.type d, sreg;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[],
            },
        ],
        "movmatrix" => &[InstructionSyntax {
            opcode: InstructionOpcode::MovMatrix,
            name: "movmatrix",
            syntax: "movmatrix.sync.aligned.shape.trans.type d, a;",
            min_operand_arity: 2usize,
            max_operand_arity: 2usize,
            required_modifiers: &[".sync", ".aligned", ".trans"],
        }],
        "mul" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::Mul,
                name: "mul",
                syntax: "mul.mode.type d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Mul,
                name: "mul",
                syntax: "mul{.rnd}.bf16 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".bf16"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Mul,
                name: "mul",
                syntax: "mul{.rnd}.bf16x2 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".bf16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Mul,
                name: "mul",
                syntax: "mul{.rnd}.f64 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".f64"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Mul,
                name: "mul",
                syntax: "mul{.rnd}{.ftz}.f32x2 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".f32x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Mul,
                name: "mul",
                syntax: "mul{.rnd}{.ftz}{.sat}.f16 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".f16"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Mul,
                name: "mul",
                syntax: "mul{.rnd}{.ftz}{.sat}.f16x2 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".f16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Mul,
                name: "mul",
                syntax: "mul{.rnd}{.ftz}{.sat}.f32 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".f32"],
            },
        ],
        "mul24" => &[InstructionSyntax {
            opcode: InstructionOpcode::Mul24,
            name: "mul24",
            syntax: "mul24.mode.type d, a, b;",
            min_operand_arity: 3usize,
            max_operand_arity: 3usize,
            required_modifiers: &[],
        }],
        "multimem.ld_reduce" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::MultimemLdReduce,
                name: "multimem.ld_reduce",
                syntax: "multimem.ld_reduce.weak{.ss}.op.type d, [a];",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".weak"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::MultimemLdReduce,
                name: "multimem.ld_reduce",
                syntax: "multimem.ld_reduce.weak{.ss}.op{.acc_prec}{.vec}.type d, [a];",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".weak"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::MultimemLdReduce,
                name: "multimem.ld_reduce",
                syntax: "multimem.ld_reduce{.ldsem}{.scope}{.ss}.op.type d, [a];",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::MultimemLdReduce,
                name: "multimem.ld_reduce",
                syntax: "multimem.ld_reduce{.ldsem}{.scope}{.ss}.op{.acc_prec}{.vec}.type d, [a];",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[],
            },
        ],
        "multimem.red" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::MultimemRed,
                name: "multimem.red",
                syntax: "multimem.red{.redsem}{.scope}{.ss}.op.type [a], b;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::MultimemRed,
                name: "multimem.red",
                syntax: "multimem.red{.redsem}{.scope}{.ss}.redop{.vec}.redtype [a], b;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[],
            },
        ],
        "multimem.st" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::MultimemSt,
                name: "multimem.st",
                syntax: "multimem.st.weak{.ss}.type [a], b;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".weak"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::MultimemSt,
                name: "multimem.st",
                syntax: "multimem.st.weak{.ss}{.vec}.type [a], b;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".weak"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::MultimemSt,
                name: "multimem.st",
                syntax: "multimem.st{.stsem}{.scope}{.ss}.type [a], b;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::MultimemSt,
                name: "multimem.st",
                syntax: "multimem.st{.stsem}{.scope}{.ss}{.vec}.type [a], b;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[],
            },
        ],
        "nanosleep" => &[InstructionSyntax {
            opcode: InstructionOpcode::Nanosleep,
            name: "nanosleep",
            syntax: "nanosleep.u32 t;",
            min_operand_arity: 1usize,
            max_operand_arity: 1usize,
            required_modifiers: &[".u32"],
        }],
        "neg" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::Neg,
                name: "neg",
                syntax: "neg.bf16 d, a;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".bf16"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Neg,
                name: "neg",
                syntax: "neg.bf16x2 d, a;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".bf16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Neg,
                name: "neg",
                syntax: "neg.f64 d, a;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".f64"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Neg,
                name: "neg",
                syntax: "neg.type d, a;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Neg,
                name: "neg",
                syntax: "neg{.ftz}.f16 d, a;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".f16"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Neg,
                name: "neg",
                syntax: "neg{.ftz}.f16x2 d, a;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".f16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Neg,
                name: "neg",
                syntax: "neg{.ftz}.f32 d, a;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".f32"],
            },
        ],
        "not" => &[InstructionSyntax {
            opcode: InstructionOpcode::Not,
            name: "not",
            syntax: "not.type d, a;",
            min_operand_arity: 2usize,
            max_operand_arity: 2usize,
            required_modifiers: &[],
        }],
        "or" => &[InstructionSyntax {
            opcode: InstructionOpcode::Or,
            name: "or",
            syntax: "or.type d, a, b;",
            min_operand_arity: 3usize,
            max_operand_arity: 3usize,
            required_modifiers: &[],
        }],
        "pmevent" => &[InstructionSyntax {
            opcode: InstructionOpcode::PmEvent,
            name: "pmevent",
            syntax: "pmevent a; pmevent.mask a;",
            min_operand_arity: 1usize,
            max_operand_arity: 1usize,
            required_modifiers: &[],
        }],
        "popc" => &[InstructionSyntax {
            opcode: InstructionOpcode::Popc,
            name: "popc",
            syntax: "popc.type d, a;",
            min_operand_arity: 2usize,
            max_operand_arity: 2usize,
            required_modifiers: &[],
        }],
        "prefetch" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::Prefetch,
                name: "prefetch",
                syntax: "prefetch{.space}.level [a]; prefetch.global.level::eviction_priority [a];",
                min_operand_arity: 1usize,
                max_operand_arity: 1usize,
                required_modifiers: &[],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Prefetch,
                name: "prefetch",
                syntax: "prefetch{.tensormap_space}.tensormap [a];",
                min_operand_arity: 1usize,
                max_operand_arity: 1usize,
                required_modifiers: &[".tensormap"],
            },
        ],
        "prefetchu" => &[InstructionSyntax {
            opcode: InstructionOpcode::Prefetchu,
            name: "prefetchu",
            syntax: "prefetchu.L1 [a];",
            min_operand_arity: 1usize,
            max_operand_arity: 1usize,
            required_modifiers: &[".L1"],
        }],
        "prmt" => &[InstructionSyntax {
            opcode: InstructionOpcode::Prmt,
            name: "prmt",
            syntax: "prmt.b32{.mode} d, a, b, c;",
            min_operand_arity: 4usize,
            max_operand_arity: 4usize,
            required_modifiers: &[".b32"],
        }],
        "rcp" => &[InstructionSyntax {
            opcode: InstructionOpcode::Rcp,
            name: "rcp",
            syntax: "rcp.approx{.ftz}.f32 d, a; rcp.rnd{.ftz}.f32 d, a; rcp.rnd.f64 d, a;",
            min_operand_arity: 3usize,
            max_operand_arity: 4usize,
            required_modifiers: &[".approx", ".f32"],
        }],
        "rcp.approx.ftz.f64" => &[InstructionSyntax {
            opcode: InstructionOpcode::RcpApproxFtzF64,
            name: "rcp.approx.ftz.f64",
            syntax: "rcp.approx.ftz.f64 d, a;",
            min_operand_arity: 2usize,
            max_operand_arity: 2usize,
            required_modifiers: &[],
        }],
        "red" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::Red,
                name: "red",
                syntax: "red{.sem}{.scope}{.space}.add.noftz{.level::cache_hint}.bf16 [a], b {, cache-policy};",
                min_operand_arity: 2usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".add", ".noftz", ".bf16"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Red,
                name: "red",
                syntax: "red{.sem}{.scope}{.space}.add.noftz{.level::cache_hint}.bf16x2 [a], b {, cache-policy};",
                min_operand_arity: 2usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".add", ".noftz", ".bf16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Red,
                name: "red",
                syntax: "red{.sem}{.scope}{.space}.add.noftz{.level::cache_hint}.f16 [a], b{, cache-policy};",
                min_operand_arity: 2usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".add", ".noftz", ".f16"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Red,
                name: "red",
                syntax: "red{.sem}{.scope}{.space}.add.noftz{.level::cache_hint}.f16x2 [a], b{, cache-policy};",
                min_operand_arity: 2usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".add", ".noftz", ".f16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Red,
                name: "red",
                syntax: "red{.sem}{.scope}{.space}.op{.level::cache_hint}.type [a], b{, cache-policy};",
                min_operand_arity: 2usize,
                max_operand_arity: 3usize,
                required_modifiers: &[],
            },
        ],
        "red.async" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::RedAsync,
                name: "red.async",
                syntax: "red.async.sem.scope{.ss}.completion_mechanism.add.type [a], b, [mbar];",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".add"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::RedAsync,
                name: "red.async",
                syntax: "red.async.sem.scope{.ss}.completion_mechanism.op.type [a], b, [mbar];",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::RedAsync,
                name: "red.async",
                syntax: "red.async{.mmio}.sem.scope{.ss}.add.type [a], b;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".add"],
            },
        ],
        "redux.sync" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::ReduxSync,
                name: "redux.sync",
                syntax: "redux.sync.op.b32 dst, src, membermask;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".b32"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::ReduxSync,
                name: "redux.sync",
                syntax: "redux.sync.op.type dst, src, membermask;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::ReduxSync,
                name: "redux.sync",
                syntax: "redux.sync.op{.abs.}{.NaN}.f32 dst, src, membermask;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".f32"],
            },
        ],
        "rem" => &[InstructionSyntax {
            opcode: InstructionOpcode::Rem,
            name: "rem",
            syntax: "rem.type d, a, b;",
            min_operand_arity: 3usize,
            max_operand_arity: 3usize,
            required_modifiers: &[],
        }],
        "ret" => &[InstructionSyntax {
            opcode: InstructionOpcode::Ret,
            name: "ret",
            syntax: "ret{.uni};",
            min_operand_arity: 0usize,
            max_operand_arity: 0usize,
            required_modifiers: &[],
        }],
        "rsqrt" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::Rsqrt,
                name: "rsqrt",
                syntax: "rsqrt.approx.f64 d, a;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".approx", ".f64"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Rsqrt,
                name: "rsqrt",
                syntax: "rsqrt.approx{.ftz}.f32 d, a;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".approx", ".f32"],
            },
        ],
        "rsqrt.approx.ftz.f64" => &[InstructionSyntax {
            opcode: InstructionOpcode::RsqrtApproxFtzF64,
            name: "rsqrt.approx.ftz.f64",
            syntax: "rsqrt.approx.ftz.f64 d, a;",
            min_operand_arity: 2usize,
            max_operand_arity: 2usize,
            required_modifiers: &[],
        }],
        "sad" => &[InstructionSyntax {
            opcode: InstructionOpcode::Sad,
            name: "sad",
            syntax: "sad.type d, a, b, c;",
            min_operand_arity: 4usize,
            max_operand_arity: 4usize,
            required_modifiers: &[],
        }],
        "selp" => &[InstructionSyntax {
            opcode: InstructionOpcode::Selp,
            name: "selp",
            syntax: "selp.type d, a, b, c;",
            min_operand_arity: 4usize,
            max_operand_arity: 4usize,
            required_modifiers: &[],
        }],
        "set" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::Set,
                name: "set",
                syntax: "set.CmpOp.BoolOp.bf16.stype d, a, b, {!}c;",
                min_operand_arity: 3usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".bf16"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Set,
                name: "set",
                syntax: "set.CmpOp.BoolOp.dtype.bf16 d, a, b, {!}c;",
                min_operand_arity: 3usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".bf16"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Set,
                name: "set",
                syntax: "set.CmpOp.BoolOp.dtype.bf16x2 d, a, b, {!}c;",
                min_operand_arity: 3usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".bf16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Set,
                name: "set",
                syntax: "set.CmpOp.BoolOp{.ftz}.dtype.f16 d, a, b, {!}c;",
                min_operand_arity: 3usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".f16"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Set,
                name: "set",
                syntax: "set.CmpOp.BoolOp{.ftz}.dtype.f16x2 d, a, b, {!}c;",
                min_operand_arity: 3usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".f16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Set,
                name: "set",
                syntax: "set.CmpOp.BoolOp{.ftz}.dtype.stype d, a, b, {!}c;",
                min_operand_arity: 3usize,
                max_operand_arity: 4usize,
                required_modifiers: &[],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Set,
                name: "set",
                syntax: "set.CmpOp.BoolOp{.ftz}.f16.stype d, a, b, {!}c;",
                min_operand_arity: 3usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".f16"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Set,
                name: "set",
                syntax: "set.CmpOp.bf16.stype d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".bf16"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Set,
                name: "set",
                syntax: "set.CmpOp.dtype.bf16 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".bf16"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Set,
                name: "set",
                syntax: "set.CmpOp.dtype.bf16x2 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".bf16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Set,
                name: "set",
                syntax: "set.CmpOp{.ftz}.dtype.f16 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".f16"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Set,
                name: "set",
                syntax: "set.CmpOp{.ftz}.dtype.f16x2 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".f16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Set,
                name: "set",
                syntax: "set.CmpOp{.ftz}.dtype.stype d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Set,
                name: "set",
                syntax: "set.CmpOp{.ftz}.f16.stype d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".f16"],
            },
        ],
        "setmaxnreg" => &[InstructionSyntax {
            opcode: InstructionOpcode::SetMaxNReg,
            name: "setmaxnreg",
            syntax: "setmaxnreg.action.sync.aligned.u32 imm-reg-count;",
            min_operand_arity: 1usize,
            max_operand_arity: 1usize,
            required_modifiers: &[".sync", ".aligned", ".u32"],
        }],
        "setp" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::Setp,
                name: "setp",
                syntax: "setp.CmpOp.BoolOp.bf16 p, a, b, {!}c;",
                min_operand_arity: 3usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".bf16"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Setp,
                name: "setp",
                syntax: "setp.CmpOp.BoolOp.bf16x2 p|q, a, b, {!}c;",
                min_operand_arity: 3usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".bf16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Setp,
                name: "setp",
                syntax: "setp.CmpOp.BoolOp{.ftz}.f16 p, a, b, {!}c;",
                min_operand_arity: 3usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".f16"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Setp,
                name: "setp",
                syntax: "setp.CmpOp.BoolOp{.ftz}.f16x2 p|q, a, b, {!}c;",
                min_operand_arity: 3usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".f16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Setp,
                name: "setp",
                syntax: "setp.CmpOp.BoolOp{.ftz}.type p[|q], a, b, {!}c;",
                min_operand_arity: 3usize,
                max_operand_arity: 4usize,
                required_modifiers: &[],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Setp,
                name: "setp",
                syntax: "setp.CmpOp.bf16 p, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".bf16"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Setp,
                name: "setp",
                syntax: "setp.CmpOp.bf16x2 p|q, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".bf16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Setp,
                name: "setp",
                syntax: "setp.CmpOp{.ftz}.f16 p, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".f16"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Setp,
                name: "setp",
                syntax: "setp.CmpOp{.ftz}.f16x2 p|q, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".f16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Setp,
                name: "setp",
                syntax: "setp.CmpOp{.ftz}.type p[|q], a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[],
            },
        ],
        "shf" => &[InstructionSyntax {
            opcode: InstructionOpcode::Shf,
            name: "shf",
            syntax: "shf.l.mode.b32 d, a, b, c; shf.r.mode.b32 d, a, b, c;",
            min_operand_arity: 7usize,
            max_operand_arity: 7usize,
            required_modifiers: &[".l", ".b32"],
        }],
        "shfl" => &[InstructionSyntax {
            opcode: InstructionOpcode::Shfl,
            name: "shfl",
            syntax: "shfl.mode.b32 d[|p], a, b, c;",
            min_operand_arity: 4usize,
            max_operand_arity: 4usize,
            required_modifiers: &[".b32"],
        }],
        "shfl.sync" => &[InstructionSyntax {
            opcode: InstructionOpcode::ShflSync,
            name: "shfl.sync",
            syntax: "shfl.sync.mode.b32 d[|p], a, b, c, membermask;",
            min_operand_arity: 5usize,
            max_operand_arity: 5usize,
            required_modifiers: &[".b32"],
        }],
        "shl" => &[InstructionSyntax {
            opcode: InstructionOpcode::Shl,
            name: "shl",
            syntax: "shl.type d, a, b;",
            min_operand_arity: 3usize,
            max_operand_arity: 3usize,
            required_modifiers: &[],
        }],
        "shr" => &[InstructionSyntax {
            opcode: InstructionOpcode::Shr,
            name: "shr",
            syntax: "shr.type d, a, b;",
            min_operand_arity: 3usize,
            max_operand_arity: 3usize,
            required_modifiers: &[],
        }],
        "sin" => &[InstructionSyntax {
            opcode: InstructionOpcode::Sin,
            name: "sin",
            syntax: "sin.approx{.ftz}.f32 d, a;",
            min_operand_arity: 2usize,
            max_operand_arity: 2usize,
            required_modifiers: &[".approx", ".f32"],
        }],
        "slct" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::Slct,
                name: "slct",
                syntax: "slct.dtype.s32 d, a, b, c;",
                min_operand_arity: 4usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".s32"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Slct,
                name: "slct",
                syntax: "slct{.ftz}.dtype.f32 d, a, b, c;",
                min_operand_arity: 4usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".f32"],
            },
        ],
        "sqrt" => &[InstructionSyntax {
            opcode: InstructionOpcode::Sqrt,
            name: "sqrt",
            syntax: "sqrt.approx{.ftz}.f32 d, a; sqrt.rnd{.ftz}.f32 d, a; sqrt.rnd.f64 d, a;",
            min_operand_arity: 3usize,
            max_operand_arity: 4usize,
            required_modifiers: &[".approx", ".f32"],
        }],
        "st" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::St,
                name: "st",
                syntax: "st.mmio.relaxed.sys{.global}.type [a], b;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".mmio", ".relaxed", ".sys"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::St,
                name: "st",
                syntax: "st.relaxed.scope{.ss}{.level1::eviction_priority}{.level2::eviction_priority}{.level::cache_hint}{.vec}.type [a], b{, cache-policy};",
                min_operand_arity: 2usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".relaxed"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::St,
                name: "st",
                syntax: "st.release.scope{.ss}{.level1::eviction_priority}{.level2::eviction_priority}{.level::cache_hint}{.vec}.type [a], b{, cache-policy};",
                min_operand_arity: 2usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".release"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::St,
                name: "st",
                syntax: "st.volatile{.ss}{.vec}.type [a], b;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".volatile"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::St,
                name: "st",
                syntax: "st{.weak}{.ss}{.cop}{.level::cache_hint}{.vec}.type [a], b{, cache-policy};",
                min_operand_arity: 2usize,
                max_operand_arity: 3usize,
                required_modifiers: &[],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::St,
                name: "st",
                syntax: "st{.weak}{.ss}{.level1::eviction_priority}{.level2::eviction_priority}{.level::cache_hint}{.vec}.type [a], b{, cache-policy};",
                min_operand_arity: 2usize,
                max_operand_arity: 3usize,
                required_modifiers: &[],
            },
        ],
        "st.async" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::StAsync,
                name: "st.async",
                syntax: "st.async{.mmio}.sem.scope{.ss}.type [a], b;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::StAsync,
                name: "st.async",
                syntax: "st.async{.scope}{.ss}.completion_mechanism{.vec}.type [a], b, [mbar];",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::StAsync,
                name: "st.async",
                syntax: "st.async{.weak}{.ss}.completion_mechanism{.vec}.type [a], b, [mbar];",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[],
            },
        ],
        "st.bulk" => &[InstructionSyntax {
            opcode: InstructionOpcode::StBulk,
            name: "st.bulk",
            syntax: "st.bulk{.weak}{.shared::cta} [a], size, initval;",
            min_operand_arity: 3usize,
            max_operand_arity: 3usize,
            required_modifiers: &[],
        }],
        "stackrestore" => &[InstructionSyntax {
            opcode: InstructionOpcode::StackRestore,
            name: "stackrestore",
            syntax: "stackrestore.type a;",
            min_operand_arity: 1usize,
            max_operand_arity: 1usize,
            required_modifiers: &[],
        }],
        "stacksave" => &[InstructionSyntax {
            opcode: InstructionOpcode::StackSave,
            name: "stacksave",
            syntax: "stacksave.type d;",
            min_operand_arity: 1usize,
            max_operand_arity: 1usize,
            required_modifiers: &[],
        }],
        "stmatrix" => &[InstructionSyntax {
            opcode: InstructionOpcode::StMatrix,
            name: "stmatrix",
            syntax: "stmatrix.sync.aligned.shape.num{.trans}{.ss}.type [p], r;",
            min_operand_arity: 2usize,
            max_operand_arity: 2usize,
            required_modifiers: &[".sync", ".aligned"],
        }],
        "sub" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::Sub,
                name: "sub",
                syntax: "sub.type1 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Sub,
                name: "sub",
                syntax: "sub{.rnd}.bf16 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".bf16"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Sub,
                name: "sub",
                syntax: "sub{.rnd}.bf16x2 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".bf16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Sub,
                name: "sub",
                syntax: "sub{.rnd}.f64 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".f64"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Sub,
                name: "sub",
                syntax: "sub{.rnd}{.ftz}.f32x2 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".f32x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Sub,
                name: "sub",
                syntax: "sub{.rnd}{.ftz}{.sat}.f16 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".f16"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Sub,
                name: "sub",
                syntax: "sub{.rnd}{.ftz}{.sat}.f16x2 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".f16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Sub,
                name: "sub",
                syntax: "sub{.rnd}{.ftz}{.sat}.f32 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".f32"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Sub,
                name: "sub",
                syntax: "sub{.rnd}{.sat}.f32.atype d, a, c;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".f32"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Sub,
                name: "sub",
                syntax: "sub{.sat}.type2 d, a, b;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[],
            },
        ],
        "sub.cc" => &[InstructionSyntax {
            opcode: InstructionOpcode::SubCc,
            name: "sub.cc",
            syntax: "sub.cc.type d, a, b;",
            min_operand_arity: 3usize,
            max_operand_arity: 3usize,
            required_modifiers: &[],
        }],
        "subc" => &[InstructionSyntax {
            opcode: InstructionOpcode::Subc,
            name: "subc",
            syntax: "subc{.cc}.type d, a, b;",
            min_operand_arity: 3usize,
            max_operand_arity: 3usize,
            required_modifiers: &[],
        }],
        "suld" => &[InstructionSyntax {
            opcode: InstructionOpcode::Suld,
            name: "suld",
            syntax: "suld.b.geom{.cop}.vec.dtype.clamp d, [a, b];",
            min_operand_arity: 3usize,
            max_operand_arity: 3usize,
            required_modifiers: &[".b"],
        }],
        "suq" => &[InstructionSyntax {
            opcode: InstructionOpcode::Suq,
            name: "suq",
            syntax: "suq.query.b32 d, [a];",
            min_operand_arity: 2usize,
            max_operand_arity: 2usize,
            required_modifiers: &[".b32"],
        }],
        "sured" => &[InstructionSyntax {
            opcode: InstructionOpcode::Sured,
            name: "sured",
            syntax: "sured.b.op.geom.ctype.clamp [a,b],c; sured.p.op.geom.ctype.clamp [a,b],c;",
            min_operand_arity: 5usize,
            max_operand_arity: 5usize,
            required_modifiers: &[".b"],
        }],
        "sust" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::Sust,
                name: "sust",
                syntax: "sust.b.{1d,2d,3d}{.cop}.vec.ctype.clamp [a, b], c; sust.p.{1d,2d,3d}.vec.b32.clamp [a, b], c;",
                min_operand_arity: 6usize,
                max_operand_arity: 7usize,
                required_modifiers: &[".b", "."],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Sust,
                name: "sust",
                syntax: "sust.b.{a1d,a2d}{.cop}.vec.ctype.clamp [a, b], c;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".b", "."],
            },
        ],
        "szext" => &[InstructionSyntax {
            opcode: InstructionOpcode::Szext,
            name: "szext",
            syntax: "szext.mode.type d, a, b;",
            min_operand_arity: 3usize,
            max_operand_arity: 3usize,
            required_modifiers: &[],
        }],
        "tanh" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::Tanh,
                name: "tanh",
                syntax: "tanh.approx.f32 d, a;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".approx", ".f32"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Tanh,
                name: "tanh",
                syntax: "tanh.approx.type d, a;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".approx"],
            },
        ],
        "tcgen05.alloc" => &[InstructionSyntax {
            opcode: InstructionOpcode::Tcgen05Alloc,
            name: "tcgen05.alloc",
            syntax: "tcgen05.alloc.cta_group.sync.aligned{.shared::cta}.b32 [dst], nCols;",
            min_operand_arity: 2usize,
            max_operand_arity: 2usize,
            required_modifiers: &[".sync", ".aligned", ".b32"],
        }],
        "tcgen05.commit" => &[InstructionSyntax {
            opcode: InstructionOpcode::Tcgen05Commit,
            name: "tcgen05.commit",
            syntax: "tcgen05.commit.cta_group.completion_mechanism{.shared::cluster}{.multicast}.b64 [mbar] {, ctaMask};",
            min_operand_arity: 1usize,
            max_operand_arity: 2usize,
            required_modifiers: &[".b64"],
        }],
        "tcgen05.cp" => &[InstructionSyntax {
            opcode: InstructionOpcode::Tcgen05Cp,
            name: "tcgen05.cp",
            syntax: "tcgen05.cp.cta_group.shape{.multicast}{.dst_fmt.src_fmt} [taddr], s-desc;",
            min_operand_arity: 2usize,
            max_operand_arity: 2usize,
            required_modifiers: &[],
        }],
        "tcgen05.dealloc" => &[InstructionSyntax {
            opcode: InstructionOpcode::Tcgen05Dealloc,
            name: "tcgen05.dealloc",
            syntax: "tcgen05.dealloc.cta_group.sync.aligned.b32 taddr, nCols;",
            min_operand_arity: 2usize,
            max_operand_arity: 2usize,
            required_modifiers: &[".sync", ".aligned", ".b32"],
        }],
        "tcgen05.fence" => &[],
        "tcgen05.ld" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::Tcgen05Ld,
                name: "tcgen05.ld",
                syntax: "tcgen05.ld.red.sync.aligned.shape3.num.redOp.type r, redval, [taddr];",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".red", ".sync", ".aligned"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Tcgen05Ld,
                name: "tcgen05.ld",
                syntax: "tcgen05.ld.red.sync.aligned.shape3.num.redOp{.abs}{.NaN}.f32 r, redval, [taddr];",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".red", ".sync", ".aligned", ".f32"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Tcgen05Ld,
                name: "tcgen05.ld",
                syntax: "tcgen05.ld.red.sync.aligned.shape4.num.redOp.type r, redval, [taddr], immHalfSplitoff;",
                min_operand_arity: 4usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".red", ".sync", ".aligned"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Tcgen05Ld,
                name: "tcgen05.ld",
                syntax: "tcgen05.ld.red.sync.aligned.shape4.num.redOp{.abs}{.NaN}.f32 r, redval, [taddr], immHalfSplitoff;",
                min_operand_arity: 4usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".red", ".sync", ".aligned", ".f32"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Tcgen05Ld,
                name: "tcgen05.ld",
                syntax: "tcgen05.ld.sync.aligned.shape1.num{.pack}.b32 r, [taddr];",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".sync", ".aligned", ".b32"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Tcgen05Ld,
                name: "tcgen05.ld",
                syntax: "tcgen05.ld.sync.aligned.shape2.num{.pack}.b32 r, [taddr], immHalfSplitoff;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".sync", ".aligned", ".b32"],
            },
        ],
        "tcgen05.relinquish_alloc_permit" => &[InstructionSyntax {
            opcode: InstructionOpcode::Tcgen05RelinquishAllocPermit,
            name: "tcgen05.relinquish_alloc_permit",
            syntax: "tcgen05.relinquish_alloc_permit.cta_group.sync.aligned;",
            min_operand_arity: 0usize,
            max_operand_arity: 0usize,
            required_modifiers: &[".sync", ".aligned"],
        }],
        "tcgen05.shift" => &[InstructionSyntax {
            opcode: InstructionOpcode::Tcgen05Shift,
            name: "tcgen05.shift",
            syntax: "tcgen05.shift.cta_group.down [taddr];",
            min_operand_arity: 1usize,
            max_operand_arity: 1usize,
            required_modifiers: &[".down"],
        }],
        "tcgen05.st" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::Tcgen05St,
                name: "tcgen05.st",
                syntax: "tcgen05.st.sync.aligned.shape1.num{.unpack}.b32 [taddr], r;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".sync", ".aligned", ".b32"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Tcgen05St,
                name: "tcgen05.st",
                syntax: "tcgen05.st.sync.aligned.shape2.num{.unpack}.b32 [taddr], immHalfSplitoff, r;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".sync", ".aligned", ".b32"],
            },
        ],
        "tcgen05.wait" => &[],
        "tensormap.cp_fenceproxy" => &[InstructionSyntax {
            opcode: InstructionOpcode::TensorMapCpFenceProxy,
            name: "tensormap.cp_fenceproxy",
            syntax: "tensormap.cp_fenceproxy.cp_qualifiers.fence_qualifiers.sync.aligned [dst], [src], size;",
            min_operand_arity: 3usize,
            max_operand_arity: 3usize,
            required_modifiers: &[".sync", ".aligned"],
        }],
        "tensormap.replace" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::TensormapReplace,
                name: "tensormap.replace",
                syntax: "tensormap.replace.mode.field1{.ss}.b1024.type [addr], new_val;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".b1024"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::TensormapReplace,
                name: "tensormap.replace",
                syntax: "tensormap.replace.mode.field2{.ss}.b1024.type [addr], ord, new_val;",
                min_operand_arity: 3usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".b1024"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::TensormapReplace,
                name: "tensormap.replace",
                syntax: "tensormap.replace.mode.field3{.ss}.b1024.type [addr], new_val;",
                min_operand_arity: 2usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".b1024"],
            },
        ],
        "testp" => &[InstructionSyntax {
            opcode: InstructionOpcode::Testp,
            name: "testp",
            syntax: "testp.op.type p, a;",
            min_operand_arity: 2usize,
            max_operand_arity: 2usize,
            required_modifiers: &[],
        }],
        "tex" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::Tex,
                name: "tex",
                syntax: "tex.base.geom.v2.f16x2.ctype d[|p], [a, {b,} c] {, e} {, f};",
                min_operand_arity: 3usize,
                max_operand_arity: 6usize,
                required_modifiers: &[".base", ".v2", ".f16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Tex,
                name: "tex",
                syntax: "tex.base.geom.v4.dtype.ctype d[|p], [a, {b,} c] {, e} {, f};",
                min_operand_arity: 3usize,
                max_operand_arity: 6usize,
                required_modifiers: &[".base", ".v4"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Tex,
                name: "tex",
                syntax: "tex.geom.v2.f16x2.ctype d[|p], [a, b, c] {, e} {, f};",
                min_operand_arity: 4usize,
                max_operand_arity: 6usize,
                required_modifiers: &[".v2", ".f16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Tex,
                name: "tex",
                syntax: "tex.geom.v2.f16x2.ctype d[|p], [a, c] {, e} {, f};",
                min_operand_arity: 3usize,
                max_operand_arity: 5usize,
                required_modifiers: &[".v2", ".f16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Tex,
                name: "tex",
                syntax: "tex.geom.v4.dtype.ctype d, [a, c] {, e} {, f};",
                min_operand_arity: 3usize,
                max_operand_arity: 5usize,
                required_modifiers: &[".v4"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Tex,
                name: "tex",
                syntax: "tex.geom.v4.dtype.ctype d[|p], [a, b, c] {, e} {, f};",
                min_operand_arity: 4usize,
                max_operand_arity: 6usize,
                required_modifiers: &[".v4"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Tex,
                name: "tex",
                syntax: "tex.grad.geom.v2.f16x2.ctype d[|p], [a, {b,} c], dPdx, dPdy {, e} {, f};",
                min_operand_arity: 5usize,
                max_operand_arity: 8usize,
                required_modifiers: &[".grad", ".v2", ".f16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Tex,
                name: "tex",
                syntax: "tex.grad.geom.v4.dtype.ctype d[|p], [a, {b,} c], dPdx, dPdy {, e} {, f};",
                min_operand_arity: 5usize,
                max_operand_arity: 8usize,
                required_modifiers: &[".grad", ".v4"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Tex,
                name: "tex",
                syntax: "tex.level.geom.v2.f16x2.ctype d[|p], [a, {b,} c], lod {, e} {, f};",
                min_operand_arity: 4usize,
                max_operand_arity: 7usize,
                required_modifiers: &[".level", ".v2", ".f16x2"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Tex,
                name: "tex",
                syntax: "tex.level.geom.v4.dtype.ctype d[|p], [a, {b,} c], lod {, e} {, f};",
                min_operand_arity: 4usize,
                max_operand_arity: 7usize,
                required_modifiers: &[".level", ".v4"],
            },
        ],
        "tld4" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::Tld4,
                name: "tld4",
                syntax: "tld4.comp.2d.v4.dtype.f32 d[|p], [a, c] {, e} {, f};",
                min_operand_arity: 3usize,
                max_operand_arity: 5usize,
                required_modifiers: &[".2d", ".v4", ".f32"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Tld4,
                name: "tld4",
                syntax: "tld4.comp.geom.v4.dtype.f32 d[|p], [a, b, c] {, e} {, f};",
                min_operand_arity: 4usize,
                max_operand_arity: 6usize,
                required_modifiers: &[".v4", ".f32"],
            },
        ],
        "trap" => &[InstructionSyntax {
            opcode: InstructionOpcode::Trap,
            name: "trap",
            syntax: "trap;",
            min_operand_arity: 0usize,
            max_operand_arity: 0usize,
            required_modifiers: &[],
        }],
        "txq" => &[InstructionSyntax {
            opcode: InstructionOpcode::Txq,
            name: "txq",
            syntax: "txq.tquery.b32 d, [a]; txq.level.tlquery.b32 d, [a], lod; txq.squery.b32 d, [a];",
            min_operand_arity: 5usize,
            max_operand_arity: 5usize,
            required_modifiers: &[".b32"],
        }],
        "vabsdiff" => &[],
        "vabsdiff2" => &[],
        "vabsdiff4" => &[],
        "vadd" => &[],
        "vadd2" => &[],
        "vadd4" => &[],
        "vavrg2" => &[],
        "vavrg4" => &[],
        "vmad" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::VMad,
                name: "vmad",
                syntax: "vmad.dtype.atype.btype.po{.sat}{.scale} d, a{.asel}, b{.bsel}, c;",
                min_operand_arity: 2usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".po"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::VMad,
                name: "vmad",
                syntax: "vmad.dtype.atype.btype{.sat}{.scale} d, {-}a{.asel}, {-}b{.bsel}, {-}c;",
                min_operand_arity: 1usize,
                max_operand_arity: 4usize,
                required_modifiers: &[],
            },
        ],
        "vmax" => &[],
        "vmax2" => &[],
        "vmax4" => &[],
        "vmin" => &[],
        "vmin2" => &[],
        "vmin4" => &[],
        "vote" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::Vote,
                name: "vote",
                syntax: "vote.ballot.b32 d, {!}a;",
                min_operand_arity: 1usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".ballot", ".b32"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::Vote,
                name: "vote",
                syntax: "vote.mode.pred d, {!}a;",
                min_operand_arity: 1usize,
                max_operand_arity: 2usize,
                required_modifiers: &[".pred"],
            },
        ],
        "vote.sync" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::VoteSync,
                name: "vote.sync",
                syntax: "vote.sync.ballot.b32 d, {!}a, membermask;",
                min_operand_arity: 2usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".ballot", ".b32"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::VoteSync,
                name: "vote.sync",
                syntax: "vote.sync.mode.pred d, {!}a, membermask;",
                min_operand_arity: 2usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".pred"],
            },
        ],
        "vset" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::VSet,
                name: "vset",
                syntax: "vset.atype.btype.cmp d, a{.asel}, b{.bsel};",
                min_operand_arity: 1usize,
                max_operand_arity: 3usize,
                required_modifiers: &[],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::VSet,
                name: "vset",
                syntax: "vset.atype.btype.cmp d.dsel, a{.asel}, b{.bsel}, c;",
                min_operand_arity: 2usize,
                max_operand_arity: 4usize,
                required_modifiers: &[],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::VSet,
                name: "vset",
                syntax: "vset.atype.btype.cmp.op2 d, a{.asel}, b{.bsel}, c;",
                min_operand_arity: 2usize,
                max_operand_arity: 4usize,
                required_modifiers: &[],
            },
        ],
        "vset2" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::VSet2,
                name: "vset2",
                syntax: "vset2.atype.btype.cmp d{.mask}, a{.asel}, b{.bsel}, c;",
                min_operand_arity: 1usize,
                max_operand_arity: 4usize,
                required_modifiers: &[],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::VSet2,
                name: "vset2",
                syntax: "vset2.atype.btype.cmp.add d{.mask}, a{.asel}, b{.bsel}, c;",
                min_operand_arity: 1usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".add"],
            },
        ],
        "vset4" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::VSet4,
                name: "vset4",
                syntax: "vset4.atype.btype.cmp d{.mask}, a{.asel}, b{.bsel}, c;",
                min_operand_arity: 1usize,
                max_operand_arity: 4usize,
                required_modifiers: &[],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::VSet4,
                name: "vset4",
                syntax: "vset4.atype.btype.cmp.add d{.mask}, a{.asel}, b{.bsel}, c;",
                min_operand_arity: 1usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".add"],
            },
        ],
        "vshl" => &[],
        "vshr" => &[],
        "vsub" => &[],
        "vsub2" => &[],
        "vsub4" => &[],
        "wgmma.commit_group" => &[InstructionSyntax {
            opcode: InstructionOpcode::WgmmaCommitGroup,
            name: "wgmma.commit_group",
            syntax: "wgmma.commit_group.sync.aligned;",
            min_operand_arity: 0usize,
            max_operand_arity: 0usize,
            required_modifiers: &[".sync", ".aligned"],
        }],
        "wgmma.fence" => &[InstructionSyntax {
            opcode: InstructionOpcode::WgmmaFence,
            name: "wgmma.fence",
            syntax: "wgmma.fence.sync.aligned;",
            min_operand_arity: 0usize,
            max_operand_arity: 0usize,
            required_modifiers: &[".sync", ".aligned"],
        }],
        "wgmma.mma_async" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::WgmmaMmaAsync,
                name: "wgmma.mma_async",
                syntax: "wgmma.mma_async.sync.aligned.shape.dtype.f16.f16 d, a, b-desc, scale-d, imm-scale-a, imm-scale-b, imm-trans-b;",
                min_operand_arity: 7usize,
                max_operand_arity: 7usize,
                required_modifiers: &[".sync", ".aligned", ".f16", ".f16"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::WgmmaMmaAsync,
                name: "wgmma.mma_async",
                syntax: "wgmma.mma_async.sync.aligned.shape.dtype.f16.f16 d, a-desc, b-desc, scale-d, imm-scale-a, imm-scale-b, imm-trans-a, imm-trans-b;",
                min_operand_arity: 8usize,
                max_operand_arity: 8usize,
                required_modifiers: &[".sync", ".aligned", ".f16", ".f16"],
            },
        ],
        "wgmma.mma_async.sp" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::WgmmaMmaAsyncSp,
                name: "wgmma.mma_async.sp",
                syntax: "wgmma.mma_async.sp.sync.aligned.shape.dtype.f16.f16 d, a, b-desc, sp-meta, sp-sel, scale-d, imm-scale-a, imm-scale-b, imm-trans-b;",
                min_operand_arity: 9usize,
                max_operand_arity: 9usize,
                required_modifiers: &[".sync", ".aligned", ".f16", ".f16"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::WgmmaMmaAsyncSp,
                name: "wgmma.mma_async.sp",
                syntax: "wgmma.mma_async.sp.sync.aligned.shape.dtype.f16.f16 d, a-desc, b-desc, sp-meta, sp-sel, scale-d, imm-scale-a, imm-scale-b, imm-trans-a, imm-trans-b;",
                min_operand_arity: 10usize,
                max_operand_arity: 10usize,
                required_modifiers: &[".sync", ".aligned", ".f16", ".f16"],
            },
        ],
        "wgmma.wait_group" => &[InstructionSyntax {
            opcode: InstructionOpcode::WgmmaWaitGroup,
            name: "wgmma.wait_group",
            syntax: "wgmma.wait_group.sync.aligned N;",
            min_operand_arity: 1usize,
            max_operand_arity: 1usize,
            required_modifiers: &[".sync", ".aligned"],
        }],
        "wmma.load" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::WmmaLoad,
                name: "wmma.load",
                syntax: "wmma.load.a.sync.aligned.layout.shape{.ss}.atype r, [p] {, stride};",
                min_operand_arity: 2usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".a", ".sync", ".aligned"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::WmmaLoad,
                name: "wmma.load",
                syntax: "wmma.load.b.sync.aligned.layout.shape{.ss}.btype r, [p] {, stride};",
                min_operand_arity: 2usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".b", ".sync", ".aligned"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::WmmaLoad,
                name: "wmma.load",
                syntax: "wmma.load.c.sync.aligned.layout.shape{.ss}.ctype r, [p] {, stride};",
                min_operand_arity: 2usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".c", ".sync", ".aligned"],
            },
        ],
        "wmma.mma" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::WmmaMma,
                name: "wmma.mma",
                syntax: "wmma.mma.sync.aligned.alayout.blayout.shape.dtype.ctype d, a, b, c;",
                min_operand_arity: 4usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".sync", ".aligned"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::WmmaMma,
                name: "wmma.mma",
                syntax: "wmma.mma.sync.aligned.alayout.blayout.shape.s32.atype.btype.s32{.satfinite} d, a, b, c;",
                min_operand_arity: 4usize,
                max_operand_arity: 4usize,
                required_modifiers: &[".sync", ".aligned", ".s32", ".s32"],
            },
        ],
        "wmma.store" => &[
            InstructionSyntax {
                opcode: InstructionOpcode::WmmaStore,
                name: "wmma.store",
                syntax: "wmma.store.d.sync.aligned.layout.shape{.ss}.type [p], r {, stride}",
                min_operand_arity: 2usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".d", ".sync", ".aligned"],
            },
            InstructionSyntax {
                opcode: InstructionOpcode::WmmaStore,
                name: "wmma.store",
                syntax: "wmma.store.d.sync.aligned.layout.shape{.ss}.type [p], r {, stride};",
                min_operand_arity: 2usize,
                max_operand_arity: 3usize,
                required_modifiers: &[".d", ".sync", ".aligned"],
            },
        ],
        "xor" => &[InstructionSyntax {
            opcode: InstructionOpcode::Xor,
            name: "xor",
            syntax: "xor.type d, a, b;",
            min_operand_arity: 3usize,
            max_operand_arity: 3usize,
            required_modifiers: &[],
        }],
        _ => &[],
    }
}