Op Mnemonic Master Table

Every MLIR operation mnemonic registered by — or observed in lowerings driven by — the tileiras ELF (CUDA Toolkit 13.1, SHA256 f0eb415767f403c96cbabf0817c3bcf70a50f88dfc8845fe36ebe21635fa6707). Nine dialect namespaces, ~640 first-class ops, alphabetical within each namespace. Columns: verbatim mnemonic in backticks, mnemonic length in bytes (- where the registrar uses a non-flat path that does not pass the literal length to RegisteredOperationName::insert), TypeID singleton sentinel (per-op &unk_NNNNNNN where the registrar exposes a per-op slot, range reference where the dialect uses a contiguous slab without per-op isolation), one-clause semantic, primary wiki page. Sentinel addresses use IDA-style &unk_NNNNNNN form, preserving the verbatim hexadecimal address from .bss/.data. The mnemonic length column matches the second argument passed to sub_4461CA0 (the RegisteredOperationName::insert callee). Where the glossary lists a range without a per-op slot, the entry cites the full range.

How the TypeID column is consumed

Every dispatcher in the binary reads OperationName::TypeID through one double-indirection from the Operation pointer:

/* The OperationName slot sits at fixed offset +0x30 on an mlir::Operation,
 * and the TypeID pointer sits at +0x10 of OperationName::Impl. Both offsets
 * are stable across the binary; every dyn_cast / OpInterface lookup in
 * tileiras decompiles to this same shape. */
static inline const void *operation_typeid(const void *op) {
    const void *opname_impl = *(const void *const *)((const uint8_t *)op + 0x30);
    return *(const void *const *)((const uint8_t *)opname_impl + 0x10);
}

/* Dispatching on an op is therefore one pointer-equality test per arm
 * against a sentinel address from the table below. A reimplementer who
 * wants the same dispatch performance must publish exactly one stable
 * address per op kind for pointer-equality identity. */
static inline bool op_is(const void *op, const void *sentinel) {
    return operation_typeid(op) == sentinel;
}

Pointer-identity sentinels (the dominant form in the slab columns below) are plain .bss slots; their address is the TypeID, no load of the byte is ever made. Meyers-cached sentinels (the cute interface anchors) hold the TypeID in a 64-bit qword that is filled in on first use through the mlir::TypeID::getFullName() interner. For the full sentinel-form breakdown and the address-band index, see TypeID Sentinel Table.

§1 cuda_tile.* (92 ops)

TypeID slab range 0x5785D0..0x57A8E0. Per-op TypeID slots are in this range but the registration thunk does not expose individual &unk_* isolated addresses to the surface decompilation; entries are cited via the range.

§2 nv_tileaa.* (73 ops)

Per-op TypeID slots in dense range 0x5B46D28..0x5B46F68 (8-byte stride). The slab anchors below the nv_tileas slab.

Note: enumeration follows the registrar walk in p2-C01:441-513 and yields 72 mnemonics including the queue.* and make_* decompositions; the "61 canonical ops" count cited in the dialect summary collapses make_memref / make_range / view to their corresponding make_* family count. All entries above are first-class.

§3 nv_tileas.* (58 ops)

Anchor &unk_5B44F08. RTTI nv_tile_ir::as. async.pipeline.* cluster dominates the surface area.

§4 cute.* (59 ops)

Anchor &unk_5B496B8. Hardware-independent CuTe layout algebra.

§5 cute_nvgpu.* (73 ops)

TypeID slab range 0x5B47FF8..0x5B481A8 (54 slots, 8-byte stride); remaining ops fall into per-op accessor singletons in same arena. Anchor &unk_5B482C8.

§6 cutlass.* (84 ops, 38 unique families)

Fold-record range 0x5B47490..0x5B476A0 covers the op-info blocks. Includes block_striped collectives, generic and named barriers, the pipeline state machine, the seq_bar protocol, and the tile_scheduler family (DP, static-persistent, StreamK, MODS-trace).

§7 mlir::nvgpu.* (upstream, observed in lowerings)

Upstream MLIR nvgpu dialect; statically linked into tileiras. Dialect TypeID anchor is provided by the upstream registration; per-op TypeIDs are not exposed by tileiras's own registrar. The list below enumerates every upstream nvgpu.* mnemonic observed in tileiras-driven lowerings (produced by convert-nvgpu-to-nvvm consumers and equivalent upstream dialects).

§8 NVVM.* (213 ops)

TypeID slab 0x5B8D610..0x5B8DCB8 (1704 bytes / 8 = 213 entries, 8-byte stride, dense). Dialect TypeID &unk_5B8DCC0 sits 8 bytes above the highest op slot. Walked via RegisteredOperationName::insert at sub_4461CA0 from the registrar driver sub_2EFC390. Order below is the categorical roster from p5-HH01 (within each category alphabetical where the registrar permits it; otherwise registrar walk order).

§8.1 Barriers (10)

§8.2 mbarrier (20)

§8.3 TMA / cp.async.bulk (12)

§8.4 cp.async (Ampere) (5)

§8.5 tcgen05 (Blackwell) (18)

§8.6 wgmma / wmma / mma / ldmatrix-stmatrix (12)

§8.7 shfl / vote / redux / match / elect (5)

§8.8 Convert / cvt.packfloat (11)

§8.9 read.ptx.sreg.* (73)

§8.10 cluster_launch_ctrl (7)

§8.11 Fences (14)

§8.12 dot_accum (2)

§8.13 griddep / proxy / tensormap (5)

§8.14 Misc (19)

§9 llvm-extras (upstream llvm.* ops observed in tileiras lowerings)

The MLIR llvm dialect is statically linked from upstream and registered via addOperation<> chains; tileiras does not surface a per-op &unk_* slot for these. The list below enumerates the llvm.* mnemonics emitted by tileiras-driven lowerings. Dialect TypeID anchor is &unk_5BA8F60.