Tileiras - MLIR-Based Optimizing Assembler

Tileiras is NVIDIA's CUDA TileIR optimizing assembler, shipped with CUDA 13.1 as a separate compiler binary. It consumes serialized MLIR bytecode for a tile program, lowers that program through NVIDIA tile dialects and NVPTX code generation, invokes the assembler toolchain, and writes a host relocatable object containing the compiled GPU payload.

The useful way to think about tileiras is not as a C++ compiler and not as a replacement for cudafe++. Tileiras starts after a frontend has already described the GPU work in MLIR. Its job is to make that tile-level program executable on Blackwell-family GPUs.

This wiki is written for two practical readers:

• If you use or integrate tileiras, the driver, option, bytecode, and subprocess pages explain what inputs the tool accepts, which target modes are valid, which external tools must be available, and how failures should be interpreted.
• If you are reimplementing compatible tooling, the subsystem pages describe observable contracts: bytecode structure, dialect schemas, pass ordering, scheduler invariants, lowering decisions, diagnostics, and pseudocode-level algorithms.

If you arrive with a specific question rather than wanting a topic tour, jump to Frequently Asked Questions, which maps common scenarios to a starting page.

At a Glance

What Tileiras Does

Tileiras is an optimizing assembler in the MLIR sense. It accepts an already-formed module, validates that the module uses the dialect versions it understands, runs a target-specific lowering pipeline, schedules and legalizes tile operations, emits PTX through the NVPTX backend, and delegates final machine-code assembly to ptxas.

The input is not CUDA C++, and tileiras does not perform preprocessing, C++ parsing, EDG lowering, template instantiation, or host-stub generation. Those responsibilities belong to other CUDA tools. Tileiras is the compiler for a lower-level tile IR surface.

The broad flow is:

tileiras bytecode
    -> parse builtin.module
    -> load cuda_tile, nv_tileaa, nv_tileas, cute, cute_nvgpu, cutlass
    -> lower tile program toward LLVM and NVVM
    -> run TileAS scheduling, layout, TMA, pipeline, and cluster passes
    -> run NVPTX code generation
    -> run ptxas
    -> optionally run nvdisasm -c for annotated disassembly payloads
    -> emit host ELF relocatable

Public Contract

For integration work, treat tileiras as a narrow bytecode-to-object compiler.

1. Produce MLIR bytecode for a builtin.module whose dialect tables match the CUDA 13.1 tile dialect schema.
2. Select one of the supported Blackwell-family targets.
3. Provide host, optimization, debug, line-info, output, and sanitizer options through the driver interface.
4. Ensure ptxas is available. Some configurations also require nvdisasm because the compile pipeline shells out to it.
5. Consume the produced object file, normally elf.o, as a host relocatable carrying the device payload.

The driver has a deliberately small option surface compared with nvcc or cicc: target GPU, host architecture, host OS, optimization level, line info, device debug, sanitizer mode, and output path. Most of the complexity is inside the bytecode reader and pass pipeline, not in command-line dispatch.

Compiler Model

Tileiras lowers across nine dialect layers. The early dialects preserve tile semantics; the middle dialects make layout, memory, and scheduling explicit; the late dialects bridge into NVVM and LLVM.

The central reimplementation point is that every stage has a structural contract. The bytecode reader must recognize the same dialect and operation tags. The pass manager must preserve the same invariants. The scheduler must obey the same resource and dependency model. The NVPTX lowering must emit the same param-space and memory-space conventions expected by ptxas.

End-to-End Algorithm

The top-level compiler can be modeled as this pipeline:

TileirasResult compile_tileiras(ByteBuffer input, TileirasConfig cfg) {
    validate_config(cfg);

    if (!is_tileiras_bytecode(input)) {
        if (looks_like_plain_mlir_bytecode(input))
            return error("failed to parse IR bytecode (it looks like MLIR bytecode instead)");
        return error("failed to parse IR bytecode");
    }

    MLIRContext ctx = create_context();
    register_tileiras_dialects(&ctx);

    Module module = parse_tileiras_bytecode(&ctx, input);
    verify_module_contract(module, cfg.gpu);

    PassManager pm = build_tileiras_pipeline(cfg);
    pm.run(module);

    LLVMModule llvm = lower_to_llvm_and_nvvm(module, cfg);
    PTXText ptx = emit_nvptx(llvm, cfg);
    Cubin sass = run_ptxas(ptx, cfg);

    Optional<Disassembly> disasm = none();
    if (cfg.requires_disassembly_payload)
        disasm = run_nvdisasm_c(sass, cfg);

    return assemble_host_object(sass, disasm, cfg.output_file);
}

The overview intentionally keeps this algorithm coarse. The detailed pages define the bytecode grammar, pass families, scheduler resource model, NVVM lowering, call-lowering ABI, and code-emission helpers at reimplementation depth.

Position in CUDA 13.1

In CUDA 13.1, tileiras is best understood as a sibling device compiler to cicc, not as a child of it. Both paths eventually produce PTX and rely on the same downstream assembler, but they start from different frontends:

CUDA C++ source path:
    CUDA C++ -> cudafe++ / cicc -> LLVM/NVVM -> PTX -> ptxas

TileIR path:
    MLIR bytecode -> tileiras -> LLVM/NVVM -> PTX -> ptxas

That distinction matters for debugging. If tileiras rejects a program, the failure is normally in bytecode schema, dialect verification, tile lowering, scheduling, NVVM conversion, or PTX assembly. It is not a C++ frontend failure.

How to Read This Wiki

The wiki is dense. Pick a path based on what you need.

For reimplementers — building a compatible CUDA TileIR compiler:

1. Start with Boundaries to fix tileiras's position in the CUDA toolchain.
2. Read Program Layout for the executable shape and subsystem map.
3. Read Pipeline Overview for the top-to-bottom cascade.
4. Drill into whichever subsystem you are implementing: dialects, scheduler, lowering, codegen, or NVPTX passes. Each subsystem page is a reimplementation-grade contract.

For users — running tileiras or diagnosing failures:

1. Start with Driver Overview for the public C-API and CLI surface.
2. Read CLI Options for the full driver option list.
3. Read Full Pass List by Opt Level to see which passes run at each -O level.
4. Jump to the specific pass page for whatever behavior you are investigating. The Reading Map curates ordered sequences for common subsystems.

For a guided tour — sample the writing quality and depth:

1. Modulo Scheduler and Rau — the scheduler exemplar, reimplementation depth.
2. MLIR Bytecode Format — the wire-format contract.
3. cuda_tile Overview — the public input IR.
4. Lowering Overview — the conversion cascade in one page.

For specific topics — see the Specialized Topics cluster and the Reading Map for curated reader paths through scheduler, codegen, dialect lowering, and OSS comparison.

Reference catalogs such as the function map, opcode rosters, and sentinel tables are intentionally denser. They are for lookup and audit work; the subsystem pages are the narrative documentation.

Documentation Style

Public pages describe behavior first. Internal recovery anchors, binary offsets, and raw analysis notes are treated as authoring evidence, not as the reader-facing API. When a recovered implementation detail matters for compatibility, the page names the semantic role first and gives pseudocode or a data-structure contract before any low-level identifier.

Code blocks use C-like pseudocode for algorithms and explicit tables for externally visible contracts. The goal is that a reader can both operate tileiras and build a compatible implementation without having to reverse the prose back into an algorithm.

Program Layout

Abstract

Tileiras is a single ELF executable with the entire MLIR-and-LLVM compiler statically linked inside it. Its segments are conventional - .text, .rodata, .data, .bss, plus the usual support sections - but the content of each segment is structured by subsystem boundary. The .text segment is partitioned into bands that correspond to the driver, the bytecode reader, dialect implementations, conversion patterns, the scheduler, codegen, and the LLVM and NVPTX libraries. The .rodata segment carries dialect string pools, pass and diagnostic strings, the bitcode writer's tag tables, and the XOR-3 obfuscated NVPTX mnemonic pool used by the asm printer. The .data segment carries cl::opt globals, dialect registration tables, and the encoded mnemonic pool that the printer walks at runtime. The .bss segment carries the singletons every subsystem accumulates: the StorageUniquer hash tables, the TypeID Meyers-cache slots, the registered dialect instances, and the per-thread caches reached through TLS. This page describes what lives in each band and why each band has the shape it does. Addresses are deliberately omitted; reimplementers care about the structure, not the offsets.

Identity

Tileiras is not a C++ frontend. It does not parse CUDA C++, instantiate templates, or generate host stubs. It consumes serialized MLIR, lowers it through the TileIR dialect stack, and produces PTX text that is then assembled by ptxas.

.text Bands

Code is grouped by subsystem; each band is a contiguous region containing the functions of one cohesive responsibility. The order, top to bottom, roughly tracks the runtime path through the compiler.

Each band has a stable internal shape: a small number of public entry points called by neighbouring bands, surrounded by a much larger population of pattern bodies and helper functions called only within the band. The lowering band is the heaviest; the scheduler band is the densest in algorithmic content per byte.

.rodata Bands

Read-only data is structured by purpose. The largest bands carry strings; the smaller bands carry the constant tables that drive the dispatch machinery.

Each rodata band has a single dominant access pattern: mnemonic pools are read by hash lookup, descriptor arrays are walked sequentially during pass registration, vtables are indexed by slot, and the libdevice blob is read end-to-end exactly once per compilation.

.data Bands

Writable but statically initialised data carries the global state that registration and option handling populate at startup.

The data segment is small relative to text and rodata. Most truly mutable state lives in .bss; the data segment exists primarily to give registration tables and cl::opt storage a stable load-time address.

.bss Bands

Zero-initialised state is where every subsystem keeps its singletons and lazy caches. These are the structures that grow during a compilation and are not meant to outlive the process.

Every entry in the .bss bands is logically owned by exactly one subsystem and is reset (or simply discarded) between compilations of independent modules.

Data Lifetimes

Three lifetime classes span the bands above. Knowing which lifetime a piece of data belongs to is what keeps the layout coherent across a long compilation.

Keeping these lifetimes separate matters in practice. cuda_tile operations are not meaningful after their conversion edge has run; scheduler analyses are not meaningful after lowering reaches LLVM; LLVM MachineIR is not meaningful before instruction selection; the XOR-3 mnemonic working copy is not meaningful before the asm printer is invoked.

Reimplementation Notes

A compatible implementation does not have to reproduce the executable's segment layout. It does have to reproduce the contracts that the layout exists to support: a stable static-init order for dialect, type, attribute, and pass registration; a uniquer that returns the same value for equal operands across the process; a .rodata-style discipline that keeps mnemonic pools, diagnostic strings, and printer tables immutable; and a .bss-style discipline that scopes all growth-only structures to the compilation that allocated them. The detailed pages under each subsystem describe these contracts as algorithms and invariants rather than as offsets.

For the reader-side view — file identity, the tools the wiki was produced with, and a recipe for verifying any individual wiki claim against the binary — see Binary Anatomy and RE Methodology.

Methodology

This wiki documents TileIR and tileiras from the behavior outward. The private reverse-engineering corpus is used as authoring evidence, but the public pages are written as technical documentation: what the program accepts, what it produces, which invariants it enforces, and how to reimplement compatible pieces.

The most important editorial rule is that implementation archaeology is not the API. A binary address, symbol nickname, or scan artifact may explain how a fact was discovered, but it usually does not help a reader operate tileiras or build compatible tooling. Public pages therefore prefer semantic names, data contracts, algorithms, and pseudocode.

Evidence Standard

Claims in this wiki are expected to rest on one of three kinds of evidence:

• direct behavior exposed by the driver, bytecode format, diagnostics, or emitted output,
• public or reconstructable MLIR/LLVM concepts such as dialect operations, pass contracts, verifier rules, and target attributes,
• repeated implementation evidence strong enough to describe a stable semantic contract.

When evidence is uncertain, pages should narrow the claim instead of exposing the uncertainty machinery. For example, a page should say "this pass materializes producer/consumer pipeline regions" only when that behavior is established. It should not ask the reader to reason from private scan identifiers.

Writing Rules

Public prose should answer the reader's practical questions first:

• What is this component for?
• Where does it sit in the pipeline?
• What inputs does it accept?
• What output or IR shape does it produce?
• What invariants must hold before and after it runs?
• What algorithm would a compatible reimplementation need?
• Which neighboring pages explain the surrounding system?

Code blocks should be reimplementation-grade pseudocode. They should name data structures and steps clearly, avoid private identifiers, and keep each line within 120 characters so the rendered wiki remains readable.

Reverse-Engineering Details

Low-level evidence still matters when it changes behavior. If a recovered detail affects compatibility, document the behavior it implies. Prefer this:

bool rrt_probe(const RRT *rrt, const NodeRRT *node, int start_cycle) {
    for (int i = 0; i < node->duration; ++i) {
        int row = (start_cycle + i) % rrt->initiation_interval;

        if ((rrt->rows[row] & node->rows[i]) != 0) {
            return false;
        }
    }

    return true;
}

Avoid presenting the same fact as a list of private function names or address ranges. That form is useful while doing the investigation, but it is not useful documentation for readers who only see the wiki.

Confidence and Corrections

Claims in this wiki carry confidence tags drawn from the three-tier HIGH/MED/LOW scheme defined in String-Evidence and Confidence Policy. The wiki's working convention is:

• inline tags (HIGH, MED, LOW) are allowed and encouraged next to individual claims when the evidence basis is worth signalling to a reader,
• a deliberately authored "Confidence" or "Evidence" section is acceptable on pages whose subject is genuinely uncertain or contested, but it is not required on every page,
• core prose should still be written from the behavior outward; if a claim is only LOW, prefer to omit it or rephrase it with hedging built into the sentence rather than parking a weak claim behind a tag.

If a later investigation changes a page, update the page inline and remove the stale claim rather than preserving a historical dispute in the main prose. When a contradiction surfaces between two analyses of the same construct, prefer the one with the stronger structural anchor regardless of recency, and restamp dependent pages.

For compatibility-sensitive behavior, prefer executable-looking algorithms and explicit invariant lists. Those are easier to review and easier to port than paragraphs describing how the fact was found.

Page Structure

Most subsystem pages should follow this shape:

1. A short purpose statement.
2. A pipeline-position diagram or paragraph.
3. The public data model or operation families.
4. Pseudocode for the central algorithm.
5. Invariants and verifier expectations.
6. Reimplementation checklist.
7. Cross-links to neighboring public pages.

Reference pages may be denser, but they should still avoid private raw evidence paths and address-heavy prose. If a table is too large to help a reader understand behavior, split it into a smaller public table and move the purely investigative detail out of the public wiki.

Subsystem Function Map

Abstract

Tileiras is a single executable, but it behaves as a stack of cooperating subsystems with a fairly rigid call-graph shape. This page is the function map for that stack at the role level: for each subsystem, the conceptual entry points it exposes, the callers it answers to, the callees it dispatches into, and the canonical wiki page where the mechanism is described in depth. The page is meant to answer the question "if I want to find where X happens in this compiler, where do I look?" without requiring a reader to know any internal address.

Call-Graph Shape at the Subsystem Level

The driver is the only entry. It builds a compile configuration, opens the input file, and calls into the bytecode reader. The reader returns an MLIR module rooted in cuda_tile. The driver hands that module, together with the configuration, to the pipeline builder. The pipeline builder asks the dialect registry for the operations and types it will see, asks the lowering subsystem for the conversion patterns it will use, and asks the scheduler for the analysis passes it must register. Pass execution then walks the registered passes in order: dialect-to-dialect lowering rewrites operations in place; the scheduler runs as an embedded phase that produces a ScheduleAnalysis without changing the IR; subsequent materialization passes consume that analysis to emit pipe and mutex operations. After the IR reaches the LLVM/NVVM dialects, the codegen subsystem translates it to an llvm::Module, links libdevice bitcode, runs the LLVM IR pipeline, and hands the result to the NVPTX backend. The backend produces PTX text. The driver then invokes ptxas as a subprocess, captures its object output, and writes the final host relocatable.

In short: driver -> bytecode -> dialects (via the registry) -> lowering -> scheduler (inside lowering) -> codegen -> libdevice link -> NVPTX backend -> ptxas. The MLIR infrastructure subsystem is orthogonal to that chain; every layer above calls into it for operations, types, attributes, regions, contexts, diagnostics, pattern rewriting, and pass management.

Driver

The driver owns process-level behavior: command-line parsing, target validation, CUDA tool discovery (ptxas, fatbinary), output naming, subprocess invocation via posix_spawn, and error reporting. Note: libdevice is not discovered here — it is embedded in the tileiras binary as _mlir_embedded_libdevice.

Callers: process entry. Callees: bytecode reader, pipeline builder, codegen, subprocess harness (ptxas, fatbinary). See Driver Overview, CLI Options, Subprocess Harness.

Bytecode Reader

The reader owns the on-disk TileIR format. It validates the container, reconstructs the string and attribute tables, and rebuilds MLIR operations and regions in memory.

Callers: driver open_input_module. Callees: dialect registry, MLIR infrastructure (storage uniquer, operation builder). See MLIR Bytecode Format.

Dialect Stack

The dialect subsystem registers operations, types, attributes, interfaces, verifiers, folders, and printers for every dialect the compiler understands. Lookup goes through OperationName and RegisteredOperationName, which is the bridge between a mnemonic and its implementation.

Each dialect exposes the same role-level entry points: register_operations, register_types, register_attributes, register_interfaces, and per-operation verify / fold / print / parse hooks. Callers: bytecode reader (to resolve mnemonics on load), lowering (to identify source-dialect operations), MLIR infrastructure (to build operations and unique types).

Lowering

The lowering subsystem owns dialect-to-dialect conversion. It is a set of pattern populations driven by ConversionTarget and DialectConversion. Each lowering page describes the patterns for one source-to-target hop.

Conversion edges, in order along the main path: cuda_tile -> nv_tileaa -> nv_tileas -> LLVM/NVGPU -> NVVM. See Lowering Overview for the conversion cascade.

Scheduler

The scheduler is a phase inside lowering. It does not modify operations; it computes a ScheduleAnalysis (stage, order, and resource placement per operation) that later materialization passes consume.

Callers: the TileAA-to-TileAS conversion edge, plus the warp-specialize pipelines. Callees: MLIR infrastructure (analysis manager, attribute storage). See Scheduler Overview and Modulo Scheduler.

Codegen and NVPTX Backend

Codegen is the boundary between MLIR and the LLVM toolchain. It produces an llvm::Module, links libdevice, runs the LLVM IR pipeline, then hands control to the NVPTX backend, which selects instructions and prints PTX.

Callers: pipeline driver, after lowering reaches LLVM dialect. Callees: libdevice subsystem, MLIR translation infrastructure. See Codegen Overview, NVPTX Backend Passes.

libdevice

libdevice is the bitcode that supplies math intrinsics. It is embedded as the symbol _mlir_embedded_libdevice inside the tileiras binary (not loaded from the CUDA toolkit), parsed once per compilation, and integrated by the LLVM IR pipeline.

Callers: codegen during the LLVM IR pipeline. See libdevice Overview, NVVM Reflect Mechanism.

MLIR Infrastructure

The infrastructure is the shared runtime model every other subsystem depends on.

Callers: every other subsystem in this map. See MLIR Infrastructure Overview, Storage Uniquer, Operation Layout.

How to Use This Map

Locate the behavior you want to understand by role: input parsing belongs to the driver and bytecode reader; semantic transformation belongs to lowering; placement decisions belong to the scheduler; libdevice math belongs to libdevice; PTX selection and emission belong to codegen and the NVPTX backend. The "page" link in each subsystem section is the canonical entry; child pages drill into one specific mechanism per page.

Versions and Fingerprints

This page records the version identifiers that matter for users and compatible implementations. It avoids private evidence anchors and focuses on the public compatibility contract: which CUDA release, LLVM lineage, dialect version family, and backend behavior this wiki describes.

Version Table

LLVM and MLIR Lineage

The key compatibility fact is that tileiras uses an LLVM/MLIR stack aligned with LLVM 21 development. That affects:

• MLIR bytecode reader behavior,
• operation, type, attribute, and interface mechanics,
• pass-manager and rewrite-pattern infrastructure,
• LLVM bitcode writing,
• NVVM intrinsic naming and lowering,
• NVPTX instruction selection and PTX emission.

A compatible reimplementation does not need to reproduce every linked LLVM helper. It does need to match the observable LLVM/NVVM contracts: data layout, target attributes, intrinsic lowering, kernel ABI, libdevice handling, and PTX backend expectations.

The binary also embeds device-side LLVM IR that was produced by older clang generations (16.0.0 and 7.1.0) running against the nvptx64-nvidia-gpulibs subtarget. The host LLVM-21 framework consumes that prebuilt IR through the standard bitcode reader; a reimplementation only needs to honor the producer-string and subtarget-triple shapes, not rebuild the embedded IR from source.

NVIDIA Extensions

The backend is not just stock upstream LLVM. It includes NVIDIA extensions for newer NVVM operations, Blackwell tensor-memory support, target-specific verifiers, NVVM reflection handling, parameter-space lowering, address-space specialization, and NVPTX machine-level cleanup.

The practical rule is:

Treat generic LLVM behavior as LLVM-21-era behavior.
Treat NVVM, NVPTX, TileIR, and tensor-memory behavior as NVIDIA-extended behavior.

When a page documents tcgen05, TMA, WGMMA, cluster launch control, TileAS scheduling, or CUTLASS pipeline lowering, assume NVIDIA-specific semantics unless the page explicitly names an upstream MLIR or LLVM feature.

External Dependency Surface

A reimplementation must account for every third-party or NVIDIA-internal component that crosses the binary's compatibility surface, not only the LLVM host link. The table below pins each one to a concrete integration point.

The integration points worth checking when a new CUDA release lands are concentrated in three places:

1. Bytecode envelope and AttrTag numbering — the wire-format fork from upstream MLIR.
2. _mlir_embedded_libdevice and the gpulibs subtarget triples — the device-side IR contract.
3. NVPTX AsmPrinter header + MatcherTable — the PTX emission contract.

Bytecode and Dialect Compatibility

The bytecode reader expects a TileIR-specific MLIR bytecode container. The public input dialect is cuda_tile; internal dialects such as nv_tileaa, nv_tileas, cute, cute_nvgpu, and cutlass are normally constructed by the pipeline or by frontend-specific producers.

Compatible tooling should preserve these boundaries:

• bytecode producers emit valid cuda_tile programs,
• dialect conversion lowers toward internal dialects in one direction,
• internal dialects are not treated as stable standalone file formats unless a page explicitly describes a textual debugging surface,
• target-specific dialects are verified against the selected compute capability.

Content Hashing

BLAKE3-style content hashing is used internally for IR object identity, deduplication, and caching. Equivalent IR objects receive stable identities within a compiler run, but the hashes are not a public ABI; treat them as implementation support.

Version-Sensitive Pages

Some pages are especially tied to CUDA 13.1 and the LLVM 21-era backend:

• NVVM Dialect Overview
• NVPTX Backend Passes
• Codegen Overview
• libdevice Overview
• MLIR Bytecode Format
• Position in nvcc 13.1

If a future CUDA release changes the bytecode schema, dialect roster, target defaults, or NVPTX intrinsic set, these pages should be reviewed first.

Glossary

This glossary defines the public terms used throughout the tileiras wiki. It focuses on behavior, data models, dialects, passes, and target concepts. Detailed operation rosters live in their dialect pages.

Core Tools

Dialects

Each dialect occupies one layer of the lowering pipeline. The early dialects preserve tile semantics, the middle dialects make layout and scheduling explicit, and the late dialects bridge to NVVM and LLVM.

Tile and Layout Terms

Scheduling Terms

Async Pipeline Terms

GPU Architecture Terms

PTX and SASS

Backend Terms

MLIR Infrastructure

Math and Precision

Reverse Engineering and Binary

CUDA Toolchain

Scheduler Coordination Values

Common Options and Environment

Reading Notes

Operation names are written in backticks, for example nv_tileas.async.pipeline.produce_one. Dialect names are also written in backticks. Pseudocode uses C-like syntax but is descriptive rather than ABI-exact unless the page explicitly says otherwise.

Reading Map

This page is curated reader paths. Each path is an ordered sequence of pages with a one-sentence rationale for why the next page follows. Use these when you want to answer "I want to understand X — what do I read in what order?" instead of browsing the SUMMARY.

Driver and Integration Path

For running tileiras, embedding it, or diagnosing a driver failure:

1. Driver Overview — what the binary does and which public entry points exist.
2. Main Entry — how main() builds the configuration and dispatches the four phases.
3. Program Handle — the 104-byte handle threaded through create / compile / get-output / release.
4. CLI Options — the option surface, separating user-facing flags from internal cl::opt plumbing.
5. Env Vars and Runtime Gates — environment-driven knobs that bypass the CLI.
6. Host Launch and ptxas Knobs — how the driver shells out to ptxas.
7. ptxas Handoff Protocol — the exact PTX surface ptxas accepts.
8. Position in nvcc 13.1 — where tileiras fits in the larger CUDA toolchain.

Bytecode Producer Path

For producing valid TileIR bytecode that tileiras will accept:

1. MLIR Bytecode Format — the container grammar and section layout.
2. Dialect Reader/Writer Status — which dialects have custom bytecode readers and what coverage looks like.
3. AsmPrinter Status — printer-side companion (the textual round-trip is partial).
4. cuda_tile Overview — the public input dialect.
5. cuda_tile Op Roster — every op the public surface accepts.
6. cuda_tile Types and Attrs — types and attributes those ops use.
7. cuda_tile Verifiers — what gets checked at parse time.
8. TypeID Sentinel Table — lookup table when you need the exact identity of a sentinel.

Dialect Lowering Chain

For understanding how the IR cascades from public input to LLVM:

1. cuda_tile — public tile-compute surface.
2. cuda_tile to tileaa — first conversion: introduce alias awareness.
3. nv_tileaa — alias-aware memory, tokens, queues.
4. tileaa to tileas — second conversion: make scheduling explicit.
5. nv_tileas — operational async-scheduling dialect.
6. cute — target-neutral layout algebra.
7. cute_nvgpu — NVIDIA architecture atoms (MMA, TMA, tcgen05).
8. cutlass — pipeline scheduler, sequence barriers, persistent kernels.
9. tileas to LLVM — final MLIR-side conversion.
10. cute and cute_nvgpu to LLVM — atom lowering to LLVM intrinsics.
11. nvgpu and gpu to NVVM — bridge to PTX-facing dialect.
12. Lowering Overview — top-down summary tying these conversions together.

Scheduler Deep-Dive

For understanding how TileAS turns dependence graphs into placed schedules:

1. Scheduler Overview — the two-pass GenerateSchedule / MaterializeSchedule split.
2. Schedule Constraint Attributes — the nine tileas.schedule.constraint.* attributes that drive placement.
3. Resource Constraint Builder and RRT — how per-op footprints become RRT bits.
4. Modulo Scheduler and Rau — the modulo-scheduling exemplar (read this one carefully).
5. Modulo Driver and 4-Arm OR-Chain — the four placement arms (PERMUTE / FUSE / RETRY / CBS).
6. Serial vs Cost-Based Generators — the two generator implementations and when each fires.
7. Schedule::solve and Cost Evaluators — the materialization algorithm.
8. Pipe and Mutex Value Layout — the IR-visible coordination values.
9. Buffer Assignment and Named Barriers — the 32-slot named-barrier pool and how Mutex_ values consume it.
10. Blackwell Pipeline 15-Slot Model — the target pipeline model the scheduler reasons against.

TileAS Pass Families

For the per-family pass roster running on nv_tileas IR:

1. Async/Pipeline Family — MaterializeSchedule, AUS vs AWS, agent materialization.
2. Layout and Buffer Family — layout assignment, slicing, and shared-memory handoffs.
3. TMA and Memops Family — TMA-descriptor and bulk-copy lowering.
4. CTA Cluster Family — cluster geometry, DynamicPersistent, PlanCTA, PrepareForScheduling, ResolveAgentBoundary.
5. Scheduling Glue — the small passes wiring schedule data into surrounding IR.

Codegen Deep-Dive

For the NVPTX backend that consumes the lowered LLVM IR:

1. Codegen Overview — pipeline shape from LLVM IR to PTX.
2. NVPTX Bring-up and Target Init — how the target gets registered and initialized.
3. NVPTX Subtarget and Feature Matrix — per-SM feature gating.
4. NVPTX Target Lowering, Call and Args — calling convention, parameter space, byval handling.
5. ISelDAG and MatcherTable — DAG-to-DAG instruction selection.
6. Per-SM Emission Templates — emission templates parameterised by SM tier.
7. AsmPrinter Monster and Windows — final PTX text emission.
8. tcgen05, WGMMA, mbarrier, Cluster — emission of the Blackwell-era instruction families.
9. TMA, Tensormap and cp.async.bulk — TMA-descriptor emission.
10. ldmatrix, stmatrix and Register Class Vtables — matrix-fragment movement.

NVPTX Custom Pass Family

For the NVIDIA-private passes layered onto the NVPTX backend:

1. NVPTX Backend Passes Overview — pipeline position and shared state.
2. Kernel, CDP, Inline, Pretreat — entry-side stamping and inline forcing.
3. Lower-Args, Aggr, Struct — byval lowering and parameter-space pointer materialization.
4. MemorySpaceOpt and process-restrict — concrete address-space inference and noalias scope generation.
5. Printf Lowering and vprintf — printf-to-vprintf rewrite.
6. DeadSyncElim and CommonBaseElim — barrier removal and SCEV-keyed GEP CSE.
7. Peephole MIR and Image Handles — post-ISel MIR rewriting.
8. NVVMIRVerifier — kernel-ABI invariants enforced before backend handoff.

libdevice and NVVM Reflect

For modules that link against libdevice math functions:

1. libdevice Overview — the bitcode library and what it covers.
2. NVVMReflect Mechanism — how compile-time reflect calls get resolved.
3. Intrinsic ID Switch and Name Table — __nv_* name to intrinsic ID mapping.
4. Math Pass Pipeline and Crosswalk — pass ordering around the math expansion.

MLIR Infrastructure Tour

For the MLIR-side mechanics referenced by dialect and lowering pages:

1. MLIR Infra Overview — what the infra layer covers.
2. Operation Layout — the 48+ byte Operation record and its slots.
3. StorageUniquer and Context Impl — type and attribute uniquing.
4. Pattern Vtables and Shapes — rewrite-pattern shapes and dispatch.
5. Interface Vtables — op and type interface mechanics.
6. TypeID Sentinels and Anchors — how TypeIDs are interned and addressed.
7. Container Fingerprints — recognizing MLIR container shapes in the binary.
8. Diagnostic ABI and Helpers — diagnostic emission, severity packing.
9. AsyncValue and BLAKE3 Interning — the 808-byte AsyncValue record backing Pipe_ / Mutex_.

OSS Comparison Tour

For comparing tileiras against the public cuda-tile repository:

1. OSS Comparison Overview — what the public tree covers vs what tileiras adds.
2. cuda_tile Tree Mapping — file-by-file mapping between public source and tileiras behavior.
3. .td Files Delta — TableGen differences.
4. Transforms, FuseFMA, SynthDbg — public transform passes and where they live in tileiras.

Cross-cutting Infra

For low-level mechanics referenced from multiple pages:

End-to-End Reimplementation Path

For a single linear read through every contract you must reproduce:

index
  -> binary-layout
  -> boundaries/nvcc-13-1-position
  -> pipeline/overview
  -> bytecode/mlir-bc-format
  -> dialects/cuda_tile/overview
  -> lowering/cuda-tile-to-tileaa
  -> dialects/nv_tileaa/overview
  -> lowering/tileaa-to-tileas
  -> dialects/nv_tileas/overview
  -> passes/tileas/scheduling-glue
  -> scheduler/overview
  -> scheduler/modulo-scheduler-and-rau
  -> lowering/tileas-to-llvm
  -> codegen/overview
  -> nvptx-passes/overview
  -> libdevice/overview

Then return to the detailed operation, verifier, and pass-family pages for the subsystem you are implementing.

Driver Overview

Abstract

tileiras is NVIDIA's TileIR optimizing assembler. It takes a TileIR bytecode module, lowers it through the TileIR and NVVM pipeline, emits PTX, invokes ptxas, and writes a host relocatable object. It is not a CUDA C++ front-end — no EDG, no cudafe, no host stub synthesis, no CUDA source parser lives in this tool. Those stages must already have produced the TileIR bytecode this driver consumes.

From the command line the driver behaves like a compact LLVM-style compiler:

tileiras [driver options] <tileir-bytecode>
    -> parse TileIR bytecode as an MLIR builtin.module
    -> run TileIR, NVVM, and NVPTX lowering
    -> serialize PTX
    -> assemble PTX with ptxas
    -> optionally dump SASS through nvdisasm -c
    -> write a host relocatable object, default elf.o

The public contract stays deliberately small. Users select the GPU architecture, host architecture, host OS, optimization/debug mode, optional memcheck instrumentation, CUDA toolkit root, and output file. The large pass inventory hiding behind that surface is catalogued in the Pipeline Overview and the Full Pass List by Opt Level.

What the driver does

One translation unit per process invocation. The input is a TileIR bytecode buffer (magic 7f 54 69 6c 65 49 52 00, version 13.1.x); the output is a host relocatable object the driver writes to --output-file or, by default, elf.o. Exit status is 0 on success or one of the five error codes documented in Driver Program Handle; no partial output is ever written.

The driver distinguishes TileIR bytecode from generic upstream MLIR bytecode at the magic-number level. A stream that opens with the MLIR framing prefix 06 03 80 0a 4d 4c 49 52 and the "\nMLIR" payload tag — rather than the TileIR "Tile\0" tag in the same slot — is rejected with a separate diagnostic that names MLIR bytecode explicitly, so the user can route the input to the right tool instead of guessing whether a parser failure means a corrupt file.

Validation runs before any pipeline construction. It rejects null buffers, non-TileIR bytecode, unsupported GPU names, optimization levels above 3, and --device-debug paired with any nonzero optimization level. The verbatim diagnostic strings and their error codes live in Driver CLI Options.

Supported Targets

The target set is Blackwell-oriented. A clean-room implementation should treat unsupported SM names as hard validation errors rather than silently remap them to the closest known architecture.

Driver Flow

The compile path is linear and has no user-visible subcommands:

main
  parse argv against the cl::opt registry
  read positional TileIR bytecode file
  resolve CUDA toolkit root
  validate buffer, target, optimization level
  create an MLIRContext and register dialects
  parse bytecode into builtin.module
  attach host/GPU target tuple
  build the TileIR pass pipeline for the requested optimization level
  lower to NVVM and LLVM
  serialize PTX text
  invoke ptxas with PTX passed as --input-as-string
  optionally write cubin to a temporary file and run nvdisasm -c
  write the relocatable object bytes to disk

The only external tools on the default path are CUDA toolkit binaries. ptxas receives PTX through --input-as-string and returns assembled cubin bytes on stdout. The SASS dump path writes that cubin to a temporary file, runs the configured disassembler command, and removes the temporary file when the driver created it.

Failure Model

Every failure prints a diagnostic and returns a nonzero exit status; the driver never writes a partial output file. The user-visible categories are:

Errors are terminal for the current invocation by design. The driver makes no attempt at partial output recovery after a pipeline or assembler failure — a fresh invocation with corrected input is always cheaper than guessing how much of a half-finished artifact is trustworthy.

Related pages

Driver main() Entry walks the entry-point code path in detail; Driver CLI Options catalogues every option and its validator; Driver Program Handle defines the public error-code numbering; Host Launch ABI and ptxas Knobs covers the kernel-launch metadata the driver emits into the produced object.

Driver main() Entry

Abstract

tileiras is a conventional LLVM-style compiler driver. The process entry point parses argv against an option schema registered during static initialization, reads the positional TileIR bytecode file, validates the buffer and the requested target, runs the TileIR-to-object pipeline inside a fresh MLIRContext, and writes the resulting host relocatable object — default path elf.o — to disk. The artifact is a relocatable object, not a raw PTX file or a standalone cubin.

The end-to-end contract is narrow. The input is an argv vector whose first non-option positional element is a TileIR bytecode file (magic 7f 54 69 6c 65 49 52 00, version 13.1.x, with at least the String and Func sections present). The output is either a host relocatable object written to --output-file and exit status 0, or one of the five integer error codes from Driver Program Handle with a verbatim diagnostic on stderr and no output file written. A failed compile leaves the filesystem in its prior state — partial output never happens.

No lowering happens in the outer driver frame. The entry point is pure orchestration: it owns option lifetime, error routing, file I/O, and the call sequence into the TileIR compiler proper.

Static-init Option Registration

Every command-line flag is an llvm::cl::opt<T> global constructed during C++ static initialization, long before main runs. Each global registers itself with the process-wide cl::OptionRegistry from its constructor, so by the time main reaches cl::ParseCommandLineOptions the parser already knows every flag, every alias, every default, and every cl::values(...) mapping table. Option objects live for the entire process lifetime and are not torn down by the driver; their storage belongs to the LLVM CommandLine library.

The four enum-valued flags (--gpu-name, --host-arch, --host-os, --sanitize) are template instantiations of cl::opt<cl::ValuesClass> that share a single parser shape: a cl::values(...) table that maps each accepted string to an int32 code, plus a default integer. The driver never sees the raw spelling — by the time the parser returns, each option holds an integer the rest of the compiler can switch on.

main()

main is pure orchestration. It runs the LLVM command-line parser, reads the positional bytecode file, builds an MLIRContext with the dialects the TileIR stack can parse or lower, dispatches to the compile entry, writes the resulting object bytes, and returns one of the five public error codes defined in Driver Program Handle.

The compile dispatcher is structured so every failure path returns before any artifact reaches disk. There is no partial-output mode and no rollback logic — a failed compile leaves the filesystem in its prior state, with diagnostics already on stderr.

int main(int argc, char **argv) {
    cl::ParseCommandLineOptions(argc, argv,
        "tileiras: NVIDIA (R) Cuda Tile IR optimizing assembler\n");

    if (InputFile.empty()) {
        errs() << "error: no input file provided\n";
        return DRIVER_ERR_INPUT_MISSING;          // code 1
    }

    auto buffer = MemoryBuffer::getFile(InputFile);
    if (!buffer) {
        errs() << "error: cannot read '" << InputFile
               << "': " << buffer.getError().message() << "\n";
        return DRIVER_ERR_INPUT_READ;             // code 4
    }

    int err = validate_driver_options(buffer->getMemBufferRef());
    if (err != 0)
        return err;                                // codes 2, 3

    MLIRContext ctx;
    register_tileiras_dialects(ctx);

    OwningOpRef<ModuleOp> module = parseSourceString<ModuleOp>(
        buffer->getBuffer(), &ctx);
    if (!module) {
        errs() << "error: failed to parse IR bytecode\n";
        return DRIVER_ERR_COMPILE;                 // code 5
    }

    attach_target(*module, GpuName, HostArch, HostOs);

    PassManager pm(&ctx);
    build_tileir_pipeline(pm, make_pipeline_options());
    if (failed(pm.run(*module))) {
        errs() << "error: failed to compile Tile IR program\n";
        return DRIVER_ERR_COMPILE;                 // code 5
    }

    SmallString<0> object;
    raw_svector_ostream os(object);
    if (failed(emit_relocatable_object(*module, os))) {
        errs() << "error: failed to emit relocatable object\n";
        return DRIVER_ERR_COMPILE;                 // code 5
    }

    auto write_err = writeToOutput(OutputFile.empty() ? "elf.o" : OutputFile,
                                   object);
    return write_err ? DRIVER_ERR_OUTPUT_WRITE : 0;
}

The branches above are the real branches: every diagnostic the driver can emit before reaching the pipeline maps to one of them. The compile path itself produces only the generic compile-failure diagnostic — finer-grained pipeline errors print from inside the pass that failed, through MLIR's diagnostic engine, before control returns to main.

Opt-level Dispatch

-O is an alias for --opt-level. The driver default is 3, the accepted range is 0..3, and the validated integer is copied into the pipeline options before the pass manager is built.

The embedded pass pipeline carries its own opt-level field with default 2, used when the pipeline runs outside the driver wrapper. The V2 pipeline carries a third field, v2-opt-level, defaulting to 0. These are three distinct axes — collapsing them produces silently different behavior for integrators who embed the pipeline directly.

Full device debug carries one hard invariant: it cannot coexist with a nonzero optimization level. With --device-debug set, the driver demands -O0, and the NVVM option string then carries debug-preserving options such as -g, --dont-merge-basicblocks, and --return-at-end. The validator rejects the combination rather than silently degrading an optimized build, because the user's intent — preserved control flow for a debugger — is incompatible with the transforms -O>0 would run.

Diagnostics Before the Pipeline

Two diagnostics fire from main before the compile dispatcher even constructs an MLIR context. A missing positional argument produces error: no input file provided and returns the input-missing code. A file that cannot be opened or mapped produces a diagnostic carrying both the path and the operating-system error message and returns the read-error code. The full validator — bytecode magic check, GPU support, optimization range, debug/optimization compatibility — runs immediately after the file is in memory and emits the verbatim diagnostics catalogued in Driver CLI Options.

--sanitize=memcheck is the only sanitizer selector accepted. Setting it appends the memcheck and tensor-memory access-check options to the downstream tool configuration.

MLIRContext and Dialect Registration

A single MLIRContext lives for the duration of one compile and owns every op, type, and attribute the pipeline allocates. Before bytecode parsing starts, the driver eagerly loads every dialect the TileIR stack can parse or lower:

Eager registration matters because MLIR bytecode references dialects by name from its symbol table. If a dialect is not loaded by the time the parser hits its first op, parsing fails with an unresolved-dialect diagnostic. The driver therefore loads every dialect the pipeline could need, and includes a late-registration fallback specifically for cuda_tile to cover bytecode that references the dialect through a form the eager path has not yet materialized.

After the module parses, the driver attaches the host/GPU target tuple through a target-attribute setter on builtin.module. The pass manager is then constructed against the module, and the pipeline is built with the options object derived from the parsed command line.

Teardown Semantics

Driver-owned cleanup is strictly local. File buffers, the MLIR context, the parsed module, and any output bytes go out of scope when main returns; LLVM's CommandLine library owns the cl::opt globals and tears them down through static destructors. MLIR's dialect registry and uniqued type/attribute storage are global runtime objects destroyed by their normal destructors after main exits — they are not part of the driver phase graph and not modeled as extra compile phases.

The distinction matters for a reimplementer building a long-lived embedding of the driver. The driver should free exactly the resources it owns and never manually destroy global dialect singletons owned by MLIR support code. Doing so corrupts the global registry for any subsequent compile in the same process.

Related pages

Driver Overview frames the surrounding pipeline from the user's perspective; Driver CLI Options catalogues every option main consumes, including the validator error codes referenced here; Driver Program Handle documents the public error-code numbering returned through main's exit status.

Driver Program Handle

Abstract

A single 104-byte program handle represents one TileIR translation unit in the tileiras driver. The public C-API surface exposes only an opaque pointer, but the recovered allocation is a fixed 0x68-byte block built by sub_57A480 (tileirasProgramCreate) and consumed by sub_57A8E0 (tileirasProgramCompile), sub_57A850 (tileirasProgramGetOutput), and sub_57A7C0 (tileirasProgramRelease). Every offset is reachable from the four public entry points, and the layout stays stable across the three create/compile/release call sites in the driver binary.

The handle stores the validated driver configuration, a small ownership bit, and an inline byte view that doubles as a CUDA-root pointer during early lifetime and as the compiled-output byte span after compile. Storage at +0x48 lives a two-phase life: the slot holds the resolved CUDA install root while the front end runs, then the same 16 bytes are repurposed to track the compiled output buffer once sub_57A8E0 finishes.

Public Error Codes

Every entry point in the driver's public C API returns a small integer status. Five non-zero codes are emitted across the four functions, and every diagnostic routes through sub_578D40 with a packed severity byte. The severity values 259, 260, and 2563 are the (class | (op_prefix << 8) | (trace << 9)) form documented in Diagnostic ABI and Helpers: 259 and 260 are fatal driver errors, 2563 is a user-input rejection.

Code 1 is reserved for the single allocation site at sub_44A8C20(0x68). Code 2 covers every configuration-level rejection. Code 3 is reserved for bytecode-parse failures and is the only code with a conditional suffix appended by a magic-tail heuristic. Code 4 is the shared null-handle / not-compiled rejection used by every entry point other than create. Code 5 is the compile-time failure code emitted by sub_57A8E0.

The MLIR fall-through on code 3 is a small heuristic inside sub_57A480. When the bytecode magic check fails, the function scans the input for the 8-byte ASCII tail e IR byt — the suffix of MLIR bytecode minus the leading ML — and on a match appends (it looks like MLIR bytecode instead) to the base code-3 diagnostic so the user can route their bytecode to the right tool.

tileirasProgramCreate Validation Order

sub_57A480 is an 820-byte routine that funnels every diagnostic through sub_578D40. Validation order is fixed and observable from the call sites — a caller can rely on each check happening before the next, because every early return is a separate diagnostic with a separate severity field.

int tileirasProgramCreate(TileirasProgram **out_program, const TileirasConfig *config) {
    if (out_program == NULL)                                    return 4;  // "program is null"
    if (config == NULL)                                         return 2;  // "configuration is null"
    if (config->input_buffer_data == NULL)                      return 2;  // "input buffer is null"
    if (!sub_57FF40(config->input_buffer_data,
                    config->input_buffer_size))                 return 3;  // parse probe; MLIR tail check appends suffix
    if (!is_supported_gpu(config->gpu_id))                      return 2;  // "unsupported GPU target"
    if ((uint32_t)config->opt_level > 3)                        return 2;  // "invalid optimization level"
    if (config->opt_level != 0 && config->device_debug == 1)    return 2;  // "optimized debugging is not supported, ..."
    if ((uint32_t)config->host_arch > 2)                        return 2;  // "unsupported host architecture"
    if ((uint32_t)config->host_os > 1)                          return 2;  // "unsupported host operating system"

    TileirasProgram *p = (TileirasProgram *)sub_44A8C20(0x68);
    if (p == NULL)                                              return 1;  // "failed to allocate memory for program"

    copy_config_into_handle(p, config);
    p->owning_flag = 0;
    *out_program = p;
    return 0;
}

The eight predicate gates are pure functions of the caller's argument tuple. Only after every predicate passes does sub_57A480 request the 104-byte block from the allocator. Allocation is the last possible failure point, so a successful return from tileirasProgramCreate guarantees the handle is initialized end-to-end.

tileirasConfig Layout

The configuration is a 16-byte aligned block whose first 80 bytes mirror the front of the program handle. sub_57A480 reads it through five _mm_loadu_si128 loads plus one scalar slot, pinning the layout to exactly five 16-byte rows.

typedef struct TileirasConfig {
    /*+0x00*/ const void *input_buffer_data;     // bytecode bytes
    /*+0x08*/ size_t      input_buffer_size;     // bytes in buffer
    /*+0x10*/ int32_t     gpu_id;                // 100/103/110/120/121 (sub_57A450 whitelist)
    /*+0x14*/ int32_t     opt_level;             // 0..3
    /*+0x18*/ int32_t     device_debug;          // 0 or 1
    /*+0x1C*/ int32_t     lineinfo;              // 0 or 1
    /*+0x20*/ int32_t     host_arch;             // 0=x86_64, 1=aarch64, 2=arm64ec
    /*+0x24*/ int32_t     host_os;               // 0=linux, 1=windows
    /*+0x28*/ int32_t     sanitize;              // 0=off, 1=memcheck
    /*+0x2C*/ uint32_t    pad_2C;
    /*+0x30*/ /* std::string SSO */              // output file name (driver side)
    /*+0x40*/ const char *cuda_root_ptr;         // resolved by sub_5773C0
    /*+0x48*/ size_t      cuda_root_len;
} TileirasConfig;

The fields at +0x10..+0x28 are the validated driver options. The CUDA root string at +0x40 is the resolution of the CUDA_ROOT / CUDA_HOME / CUDA_PATH environment chain (with /proc/self/exe fallback) performed by sub_5773C0; tileiras carries it into the program handle because the compile dispatch needs an installation path to invoke nvdisasm.

TileirasProgram Layout

The 104-byte program handle reuses the first 80 bytes of the configuration almost verbatim, with one deliberate field reorder: gpu_id moves to +0x28 so the validated 32-bit fields at +0x10..+0x28 form a contiguous scalar block that the compile dispatch reads via aligned 32-bit loads.

typedef struct TileirasProgram {
    /*+0x00*/ const void *input_buffer_data;     // ptr to bytecode bytes
    /*+0x08*/ size_t      input_buffer_size;     // bytes in buffer
    /*+0x10*/ int32_t     opt_level;             // 0..3
    /*+0x14*/ int32_t     device_debug;          // 0 or 1
    /*+0x18*/ int32_t     lineinfo;              // 0 or 1
    /*+0x1C*/ int32_t     host_arch;             // 0=x86_64, 1=aarch64, 2=arm64ec
    /*+0x20*/ int32_t     host_os;               // 0=linux, 1=windows
    /*+0x24*/ int32_t     sanitize;              // 0 or 1 (memcheck)
    /*+0x28*/ int32_t     gpu_id;                // 100/103/110/120/121
    /*+0x2C*/ uint32_t    pad_2C;
    /*+0x40*/ const char *cuda_root_ptr;         // resolved at compile time
    /*+0x48*/ size_t      cuda_root_len;         // ── overlapping slot ──
    /*+0x48*/ uint8_t    *output_data;           // SSO: same 16 bytes, second life
    /*+0x50*/ uint64_t    output_capacity;
    /*+0x58*/ size_t      output_length;
    /*+0x60*/ uint32_t    owning_flag;           // 1 = handle owns output_data
} TileirasProgram;

Five unaligned SSE copies from the configuration populate the block. The first copy moves input_buffer_data and input_buffer_size; the next two move the eight 32-bit option fields; the fourth clears the slot at +0x30; the fifth installs the CUDA-root SSO. owning_flag at +0x60 clears to zero before sub_57A480 returns, so the handle starts life with no output buffer to free.

SSO Overlap at +0x48

The slot at +0x48 is the only address in the handle that hosts two different fields across the program lifetime. While the front end runs, +0x40..+0x4F carries a (cuda_root_ptr, cuda_root_len) pair pointing at the resolved CUDA install path. The compile dispatcher reads the pair once, uses it to locate the nvdisasm binary, and never touches it again. The same 16 bytes are then overwritten to hold the compiled-output buffer descriptor: output_data at +0x48, output_capacity at +0x50, output_length at +0x58.

The offsets used by the consumers make this observable. sub_57A8E0 writes output_data at +72 (0x48), output_capacity at +80 (0x50), and output_length at +88 (0x58). sub_57A850 returns {ptr=+0x48, length=+0x58} to the caller. The lifetimes never overlap: by the time sub_57A8E0 installs the output bytes, the CUDA-root string has already been consumed and is no longer needed by any subsequent stage.

A reimplementation should keep this overlap as an internal storage trick and never expose it to callers. The public contract is simply that the program-output getter is invalid until compile has succeeded.

Ownership and Release

A single bit at +0x60, owning_flag, controls release behavior. It starts at 0 immediately after tileirasProgramCreate. sub_57A8E0 sets it to 1 once the compiled byte buffer has landed at +0x48..+0x58. sub_57A7C0 (tileirasProgramRelease) reads the flag before tearing down: 1 makes the release path call the buffer's deleter on output_data; 0 leaves the slot alone. Either way, the 104-byte handle itself is freed via the matching sub_4580C60 deallocator after the conditional output free.

The ownership rule means calling tileirasProgramRelease on a never-compiled handle is safe and touches no output memory. It also means the only legal way to retain compiled bytes after release is to copy them out via tileirasProgramGetOutput first.

Public-API Surface

Four C-API entry points operate on the handle. Each one is a small wrapper that validates its arguments, walks a fixed offset path, and routes diagnostics through sub_578D40.

All four entry points share the same C-API error space. The diagnostic emitter at sub_578D40 is the single sink for every public message, which is why codes overlap across functions: code 4 covers program is null, output pointer is null, and program has not been compiled from the getter and release sites, while code 5 is reserved for the compile-time failure path inside sub_57A8E0.

Lifecycle

The normal lifetime runs create → compile → get output → release, and the driver's main follows that sequence exactly. The handle is opaque at the public surface; the offset layout above is an implementation detail that reimplementers should reproduce for binary compatibility with the recovered driver but should never surface to consumers.

TileirasProgram *program = NULL;
int err = tileirasProgramCreate(&program, &config);
if (err == 0) {
    err = tileirasProgramCompile(program);
}
if (err == 0) {
    TileirasByteView out;
    err = tileirasProgramGetOutput(program, &out);
    if (err == 0) {
        write_output_file(opts.output_path, out.data, out.length);
    }
}
tileirasProgramRelease(program);

Compile is allowed to fail after a successful create. When it does, sub_57A8E0 returns code 5 (failed to compile Tile IR program), owning_flag stays 0, the output slot still holds the CUDA-root pair, and release frees only the handle. The getter rejects subsequent calls with code 4 (program has not been compiled), preserving the invariant that a returned byte view stays valid until the next mutating call on the handle.

Reimplementation Invariants

The handle is 104 bytes. The option fields at +0x10..+0x28 are eight contiguous 32-bit integers in the order (opt_level, device_debug, lineinfo, host_arch, host_os, sanitize, gpu_id, pad). The slot at +0x48..+0x57 is a single 16-byte region with two sequential lifetimes — CUDA-root SSO during the front end, output-buffer descriptor afterwards. The byte at +0x60 is the ownership flag, the only field tileirasProgramRelease inspects when deciding whether to free the output buffer. Allocation goes through sub_44A8C20(0x68) and deallocation through sub_4580C60; no intermediate reallocation occurs on the create or compile paths.

Cross-References

Driver main() Entry documents how the handle is threaded through the four-phase driver and how the configuration is built from cl::opt storage. The compile dispatcher that mutates the handle is described in the TileIR Pipeline Overview.

Driver CLI Options

Abstract

The tileiras command-line surface has two layers. Normal users see only the first — a small driver layer with input file, output file, target selection, optimization level, debug mode, line info, and memcheck. The second is the TileIR pipeline option structure, which surfaces when the driver constructs the pass pipeline or when an integrator embeds the pipeline directly.

The two layers reuse a few names on purpose. Driver --opt-level defaults to 3; the embedded pipeline option named opt-level defaults to 2. Treat them as separate axes unless the driver has explicitly copied the command-line choice into the pipeline options.

Driver Options

The driver parses these with LLVM command-line semantics — aliases are exact aliases, boolean flags follow LLVM's normal spelling rules, and unknown options are rejected before any compilation work starts.

Enum-valued Options as int32 Codes

The four enum-valued driver options — --gpu-name, --host-arch, --host-os, --sanitize — are wired through cl::opt<cl::ValuesClass> template instantiations that share one parser shape. Each option carries its own cl::values(...) mapping table that pairs an accepted spelling with an int32 code, plus a default integer to use when the option is absent. The parser walks the table once at command-line time, stores the resulting integer, and downstream code never sees the original string.

--gpu-name maps spellings to the corresponding SM number and defaults to 100:

The driver surface accepts only the bare sm_NN spelling — a and f variant suffixes are not parsed here. The architecture-specific selection happens one level up, on the nv_tileaa.compute_capability module attribute set by the frontend. A frontend that lowers WGMMA, tcgen05.mma, or block-scaled mma.sync carries the matching target_spec field on the module; the backend reads both fields when constructing the NVPTX target machine, picks sm_100a (for example) instead of sm_100, and emits .target sm_100a accordingly. --gpu-name is therefore a defaulting hint for the major SM number, not the final word on the .target line. The full subtarget-construction mechanism — including how --gpu-name combines with +ptxNN feature flags to drive the .version/.target header — is documented in PTX Version and Target Selection.

Two practical consequences follow. First, a kernel emitted by a frontend that requires arch-conditional instructions cannot be redirected to a plain sm_NN target by changing --gpu-name alone — the lowering will fail in the selector when no legal MachineInstr is found. Second, this driver does not list sm_90: its primary deployment surface is Blackwell, and Hopper targets are reachable only through the frontend's own attribute writes plus a host environment that pins the build to an sm_90-capable subtarget table.

--host-arch defaults to 0:

--host-os defaults to 0:

--sanitize defaults to 0 and is the only option whose unset state carries semantic weight downstream:

The host-architecture lookup table is keyed by code and walked with two strides — 39 for the x86_64 record and 36 for both aarch64 entries. arm64ec reuses the aarch64 record at a distinct sub-entry; that sub-entry is the only place the two ARM modes diverge in the host-side code path. The host-OS index resolves to 7 for linux and 15 for windows, both of which select a target-triple OS fragment and the matching object-file format.

Each parser exposes an 8-slot vtable shared by all four options. The slots are: typeinfo helper, destructor, parse (string → int32 map probe), print (int32 → string lookup for --help), valuesDefault initialiser, and three reserved slots. parse is the only operation invoked at command-line time; print fires only when the user requests help text.

Validation Algorithm

The option validator is deliberately strict. It checks the bytecode buffer and the requested target before allocating the program handle, keeping failure paths simple and steering clear of partially initialized session state.

int validate_driver_options(const ByteSpan *input, const DriverOptions *opts) {
    if (input == NULL || input->data == NULL)
        return error("input buffer is null");                                     // code 2

    if (!is_tileir_bytecode(*input)) {
        if (looks_like_mlir_bytecode(*input))
            return error("failed to parse IR bytecode (it looks like MLIR bytecode instead)");  // code 3
        return error("input does not correspond to Tile IR bytecode");            // code 3
    }

    if (!is_supported_gpu(opts->gpu_name))
        return error("unsupported GPU target");                                   // code 2

    if ((uint32_t)opts->opt_level > 3)
        return error("invalid optimization level");                               // code 2

    if (opts->device_debug && opts->opt_level != 0)
        return error("optimized debugging is not supported, "
                     "change optimization level to 0 or disable full debug info"); // code 2

    return 0;
}

The diagnostic strings above are the verbatim messages emitted by the validator entry point; the full error-code table with severity bytes lives in Driver Program Handle. The debug rule is not cosmetic — full device debug mode injects NVVM debug options that disable several code-motion and block-merge transforms, so the driver demands -O0 rather than silently degrading an optimized build.

Pipeline Options

The TileIR pass pipeline carries a much larger option structure. These options matter most to integrators who build a pass pipeline directly or expose advanced tuning flags in a higher-level tool.

The two scheduler knobs — rrt-size-threshold and max-constraint-iterations — are compile-time controls. Lower thresholds compress the resource reservation table earlier; lower iteration caps make the solver stop sooner and fall back to conservative scheduling when constraints remain unresolved.

Effective Option Merge

A TileIRPipelineOptions value is the resolved configuration that reaches the pass manager. The driver builds it in three layers, applied in order; each layer can only overwrite fields the next layer explicitly touches, so the precedence is unambiguous.

The first layer is the TableGen-declared per-field default. Every option in the pipeline has a default literal written into the pass definition — opt-level = 2, num-warps = 4, rrt-size-threshold = 4096, and so on. Constructing a fresh TileIRPipelineOptions populates every field with this baseline.

The second layer is the per-pass override that arrives through MLIR's --pass-pipeline="tileir{key=value, ...}" syntax. When the user (or an integrator) invokes the pipeline through that surface, MLIR's option parser walks the brace-enclosed key=value list and writes each value into the matching pipeline field, leaving every other field at its TableGen default.

The third layer is driver-level legacy propagation. The command-line driver predates the per-pass options syntax, and several user-facing flags — --opt-level, --gpu-name, --lineinfo, --device-debug, --sanitize, --host-arch, --host-os — must continue to work for users who never type a --pass-pipeline argument. The driver therefore copies each of those into the corresponding pipeline field after the first two layers have settled.

TileIRPipelineOptions make_pipeline_options(const DriverFlags &flags) {
    TileIRPipelineOptions opts;                           // TableGen defaults

    if (flags.pass_pipeline_set)
        parsePassPipelineOptions(opts, flags.pass_pipeline_text);  // brace-list overrides

    opts.opt_level          = flags.opt_level;            // legacy propagation
    opts.compute_capability = sm_number_of(flags.gpu_name);
    opts.emit_line_info     = flags.lineinfo ? LineInfo::FromInput : LineInfo::None;
    opts.device_debug       = flags.device_debug;
    opts.sanitize_memcheck  = flags.sanitize == Sanitizer::Memcheck;
    opts.host_arch          = flags.host_arch;
    opts.host_os            = flags.host_os;
    return opts;
}

The propagation exists because a single --opt-level=2 should still configure the pipeline correctly without forcing the user to spell out --pass-pipeline="tileir{opt-level=2}". A reimplementer who skips the propagation step ends up with a tool whose -O2 silently runs at the pipeline default of 2 for most fields but at the driver default of 3 in any field the driver does not propagate — a subtle divergence that turns up only when an integrator's regression suite compares produced SASS across versions.

Do not collapse v2-opt-level into driver --opt-level. The two are independent axes: v2-opt-level defaults to 0 and is only meaningful inside the V2 pipeline, which the driver does not select on its own.

Diagnostics Surface

Four options produce artifacts useful for debugging:

The driver does no semantic check on these paths beyond ordinary file I/O. When a path is set, the corresponding pipeline stage owns the write and reports failure through the normal compile error path.

Related pages

Driver main() Entry shows how main consumes the parsed options; Driver Overview frames the overall compile contract; Driver Program Handle defines the public error-code numbering returned through the exit status; Host Launch ABI and ptxas Knobs covers --knobs-file=, the only ptxas-side option the driver forwards. Debugging and Introspection is the user-facing debugging surface: it ties the four diagnostic options in the table above (--lineinfo, emit-line-info, schedule-trace-file, dump-host) to the MLIR-side print, timing, and stack-trace flags and gives a symptom-to-flag decision matrix. Troubleshooting and Known Issues catalogs the verbatim rejection strings produced by the validator above (unsupported GPU target, invalid optimization level, optimized debugging is not supported, could not find libdevice), pairs each with its root cause, and lists the gotchas that the strict CLI parser surfaces — notably that --gpu-name does not accept the a/f arch-conditional suffix and that sm_90 is not in the accept table.

Driver Env Vars + Runtime Gates

Abstract

tileiras carries two configuration surfaces beyond ordinary command-line options. First, a small set of environment variables drives toolkit discovery, ptxas knob forwarding, TMA policy, swizzle selection, and a TileAS shared-memory debug escape hatch. Second, a family of pass-level runtime gates lives as LLVM command-line options — not process environment variables, but worth documenting here because they shape the same compile-time behavior.

Environment-variable parsing is uneven by design. Some variables are presence-only, some require an exact value, and one parses a base-10 integer. A faithful reimplementation should preserve those differences rather than collapse every variable into a generic boolean parser.

Per-Feature Runtime Gates

Each variable has a known consumer sub_ADDR and a known parse mode. The parse modes vary on purpose: some are presence-only, one demands the literal byte "1", one parses a base-10 integer through strtol, and one is a path forwarded verbatim into an argv slot. A faithful reimplementation must preserve those differences rather than collapse them into a single boolean parser.

Three operational details matter. First, MLIR_ENABLE_EVO and PTX_KNOBS_PATH form an AND gate inside the ptxas-argv builder — setting only one does not forward a knob file. Second, the presence-only gates do not strip the value: assigning "0" or "false" still enables them, opposite of what a deployment-script reader usually expects. Third, TILE_AS_DEBUG_VERBOSE is a string-equality compare against "1", so "true" and "yes" are silently ignored.

The TILEIR_DELAY_TMA_STORE_WAIT hazard deserves its own paragraph because the failure mode is brutal. The consumer parses the value with strtol, but the calling path treats parse failure as an exception (a std::stoi-shaped std::invalid_argument) that percolates up unhandled. The result is SIGABRT on terminate. Setting TILEIR_DELAY_TMA_STORE_WAIT=foo instead of TILEIR_DELAY_TMA_STORE_WAIT=0 produces no warning and no fallback to default — it aborts the compile. Deployment scripts should either omit the variable entirely or set a decimal integer.

CUDA-Root Resolution

CUDA-root resolution is trickier than it looks: two separate resolvers live in the binary and they disagree on the miss path. The driver-side resolver sub_5773C0 fires early during command-line processing; the NVVM-side resolver sub_1A41D30 fires later from inside libdevice and libnvvm lookup. Both walk the same env-var chain — CUDA_ROOT, then CUDA_HOME, then CUDA_PATH — but they part ways when every variable is unset.

The divergence is the hazard. A deployment that leaves CUDA_ROOT unset but keeps /proc/self/exe resolvable inside the expected toolkit layout sees the driver succeed silently. Later, when the NVVM path tries to locate libdevice.10.bc, its resolver returns "", and the libdevice loader joins that empty string against nvvm/libdevice/libdevice.10.bc. The user gets a confusing "libdevice.10.bc not found in $CUDA_ROOT/nvvm/libdevice" error even though the driver itself reported no problem. The workaround is mechanical: always export CUDA_ROOT explicitly in production deployments, even when /proc/self/exe would have been enough for the driver alone.

Runtime-gate globals (static-ctor populated)

The runtime-gate layer is ordinary LLVM option storage for individual passes. These flags help when debugging pass behavior, but getenv never reads them.

Consumers

Each environment variable is consumed close to the operation it affects:

PtxasArgs build_ptxas_args(const DriverState *state) {
    PtxasArgs args = default_ptxas_args(state);

    if (getenv("MLIR_ENABLE_EVO") != NULL) {
        const char *knobs = getenv("PTX_KNOBS_PATH");
        if (knobs != NULL)
            args_append(&args, concat("--knobs-file=", knobs));
    }

    return args;
}

bool unlimited_smem_debug(void) {
    return getenv("TILE_AS_DEBUG_UNLIMITED_SMEM") != NULL;
}

bool verbose_debug(void) {
    const char *value = getenv("TILE_AS_DEBUG_VERBOSE");
    return value != NULL && strcmp(value, "1") == 0;
}

int delay_tma_store_wait(void) {
    const char *value = getenv("TILEIR_DELAY_TMA_STORE_WAIT");
    if (value == NULL)
        return 0;
    // Parse path raises std::invalid_argument on bad input;
    // calling frame does not catch, so the process aborts.
    return std::stoi(value);
}

The two CUDA-root resolvers look structurally similar, but their miss paths are not equivalent. The pseudocode below is the contract a reimplementation must preserve byte-for-byte — in particular, the NVVM resolver must keep its empty-string return, because downstream libdevice loading relies on that sentinel to defer the actual lookup to a layered fallback that lives outside this function.

const char *resolveCudaRoot_driver(void) {              // sub_5773C0
    if (const char *p = getenv("CUDA_ROOT"))   return p;
    if (const char *p = getenv("CUDA_HOME"))   return p;
    if (const char *p = getenv("CUDA_PATH"))   return p;
    return walkSelfExeUpTwo();                          // /proc/self/exe -> ../../
}

const char *resolveCudaRoot_nvvm(void) {                // sub_1A41D30
    if (const char *p = getenv("CUDA_ROOT"))   return p;
    if (const char *p = getenv("CUDA_HOME"))   return p;
    if (const char *p = getenv("CUDA_PATH"))   return p;
    return "";                                          // hazard: no fallback
}

Related pages

Host Launch ABI and ptxas Knobs covers how the PTX_KNOBS_PATH value lands in the ptxas argv; Subprocess Harness shows the ptxas-launcher argv shape that consumes the forwarded knob path; Driver CLI Options documents the companion option flags whose semantics overlap with these gates; Resource Constraint Builder and RRT is the RCB consumer of TILE_AS_DEBUG_UNLIMITED_SMEM.

Host Launch ABI + ptxas Knobs

Abstract

tileiras is assembler-side. It never calls cuLaunchKernel, cuLaunchKernelEx, or cuKernelSetAttribute directly — instead it emits kernel launch metadata into IR attributes and PTX directives, and ptxas lifts that information into the cubin metadata consumed by the CUDA runtime or driver.

The host-visible launch ABI splits across three channels:

1. PTX directives in each kernel .entry header.
2. MLIR nvvm.* attributes on lowered LLVM functions.
3. gpu.launch_func and nv_tileaa.launch_func properties that carry dynamic launch operands through the lowering pipeline.

The ptxas --knobs-file=<path> path is separate. tileiras forwards the argument only when the environment gate is enabled; ptxas owns the file grammar and every diagnostic.

Host-side launch ABI

Since the driver never synthesizes CUDA-driver launch calls, the compiled cubin carries static launch metadata and leaves dynamic launch assembly to the consumer. Static metadata flows from nvvm.* attributes and PTX directives; dynamic metadata rides on launch-operation properties and segment-size arrays during MLIR lowering.

The split that matters:

Cluster directives are gated to SM90 and newer. On older targets the compiler suppresses .blocksareclusters, .explicitcluster, .reqnctapercluster, and .maxclusterrank even when cluster-shaped metadata is present upstream.

gpu.launch_func carries kernelFunc, kernelModule, and operandSegmentSizes. The setter also accepts the older operand_segment_sizes spelling for compatibility with MLIR v17-era IR. By the nv_tileaa.launch_func stage the kernel reference flattens into a single kernel property alongside the same operand segment sizing.

nvvm.* Annotations and PTX Directives

The nvvm.* attribute family is the canonical in-IR carrier of launch metadata. Legacy !nvvm.annotations tuples still parse and can be transplanted into attribute form; an internal marker prevents repeated legacy scans after the transplant.

The verifier enforces the shape rules that matter:

1. Dimensional attributes contain one to three i32 values, except cluster dimensions, which require three values.
2. Scalar resource attributes are integer attributes.
3. nvvm.blocksareclusters requires both nvvm.reqntid and nvvm.cluster_dim on the same function.

Several invariants are core for reimplementers. nvvm.maxclusterrank is stored as an integer-valued function attribute, unlike the string-shaped legacy forms used by some older launch metadata. local_maxnreg has no new nvvm.* mirror — it stays legacy-only and is never printed as a PTX directive by this stage. When updating dimensional attributes, write every axis back together so the new attribute form stays coherent even when the legacy source used split per-axis tuples.

PTX emission walks the verified attribute set in a fixed order so that related directives stay adjacent in the kernel header. The thread-shape group (.maxntid, .reqntid) emits first, followed by the residency hints (.minnctapersm, .maxnreg), and finally the cluster group (.blocksareclusters, .explicitcluster, .reqnctapercluster, .maxclusterrank) when the target supports clusters. Both .maxntid and .reqntid may appear on the same kernel — the PTX semantics make them complementary: .reqntid declares an exact block shape the kernel relies on, .maxntid declares an upper bound for register-pressure budgeting. The verifier checks shape consistency but does not collapse or override either directive, and the emitter prints both as written when both are set.

void emit_launch_directives(LLVMFuncOp fn, Target target, PTXWriter &out) {
    if (auto dims = get_dim_attr(fn, "nvvm.maxntid"))
        out.directive(".maxntid", *dims);
    if (auto dims = get_dim_attr(fn, "nvvm.reqntid"))
        out.directive(".reqntid", *dims);

    if (auto n = get_int_attr(fn, "nvvm.minctasm"))
        out.directive(".minnctapersm", *n);
    if (auto n = get_int_attr(fn, "nvvm.maxnreg"))
        out.directive(".maxnreg", *n);

    if (!target_supports_clusters(target))
        return;                              // suppress all cluster directives pre-SM90

    if (fn->hasAttr("nvvm.blocksareclusters"))
        out.directive(".blocksareclusters");
    if (auto dims = get_dim_attr(fn, "nvvm.cluster_dim")) {
        out.directive(".explicitcluster");
        out.directive(".reqnctapercluster", *dims);
    }
    if (auto n = get_int_attr(fn, "nvvm.maxclusterrank"))
        out.directive(".maxclusterrank", *n);
}

Two structural invariants keep this loop from being more complex. nvvm.blocksareclusters is verified to require both nvvm.reqntid and nvvm.cluster_dim on the same function, so by the time emission runs the three directives are guaranteed to form a coherent triple. Cluster directives are suppressed wholesale on pre-SM90 targets; the verifier permits the attributes upstream so a single IR module can lower for multiple targets, but the per-target emitter refuses to print them when ptxas would reject the result.

How tileiras chooses each directive

Verifying an attribute is well-formed is not the same as choosing its value. The well-formedness rules above guard against malformed PTX; the choice of value is what determines whether the kernel runs at all and how fast it runs when it does. Each directive has its own input channel — the kernel-spec attribute the upstream lowering attaches, a user-supplied annotation that survives the front-end, or a constraint imposed by an instruction the compiler emitted later. The table below walks the policy for each directive.

The .maxntid versus .reqntid distinction is the policy decision that affects the most kernels. .maxntid is an upper bound the launch must respect but does not have to saturate; the driver accepts any launch shape with X, Y, Z components no larger than the declared maxima. .reqntid is a hard contract — the driver rejects any launch whose block shape does not match the declared values exactly. Tileiras emits .reqntid whenever the lowering has already baked in a specific thread count: any kernel that emits WGMMA needs 128 threads per CTA (four warps form one warp group), any kernel with warp-specialized producer/consumer splits needs the exact named warp count, and any kernel with named-warp NamedBarrier slots needs the exact thread count the slot binding assumed. For kernels that adapt to launch shape — elementwise kernels, kernels that use only synchronous mma.sync forms, kernels with no warp specialization — tileiras emits only .maxntid so the same cubin works for a range of launch shapes.

The .maxnreg choice is similarly central to performance. A WGMMA-using kernel must leave room for the accumulator fragment: an m64n256k16 FP32 WGMMA needs 32 FP32 registers per thread just for the accumulator, plus the working set for descriptors, loop indices, and any other live values. Setting .maxnreg too low forces ptxas to spill the accumulator to local memory, which silently regresses throughput by an order of magnitude. Setting it too high reduces occupancy and hurts latency hiding. The kernel-spec carries the result of a per-kernel computation that balances both — usually accumulator_regs + descriptor_regs + slack, with slack calibrated to the SM's register file size and the desired CTAs per SM.

Cluster-shape directives are an all-or-nothing group. When the kernel-spec carries nvvm.cluster_dim, the lowering emits .explicitcluster, .reqnctapercluster, and any nvvm.maxclusterrank or nvvm.blocksareclusters markers; when the spec is silent, no cluster directive is emitted. The verifier rule that nvvm.blocksareclusters requires both nvvm.reqntid and nvvm.cluster_dim means the three-directive triple is always coherent by the time the emitter sees it.

GPU Execution Model is the canonical reference for how the five tiers (thread, warp, CTA, cluster, grid) consume these directives at runtime, with a worked example that traces the directive emission from kernel-spec to PTX header for a Hopper GEMM.

ptxas Knobs File Format

When both MLIR_ENABLE_EVO and PTX_KNOBS_PATH are set, tileiras forwards --knobs-file=<path> to ptxas. It does not parse or validate the file — the grammar belongs to ptxas.

The file format is:

arbitrary preamble
[knobs]
command command command

The [knobs] sentinel is case-sensitive; text before it is ignored. After the sentinel, whitespace, ~, and ;; separate commands. The command stream has no quoting, no escaping, no comment syntax.

Commands have three forms:

Values parse per the knob descriptor type. The recovered parser accepts signed and unsigned integers, integer ranges, integer lists, 32-bit and 64-bit floats, strings, pointers, opcode lists, opcode-pair lists, and WHEN clauses. Integer parsing is decimal — a string like 0x10 parses as zero, with the trailing text ignored by the numeric conversion path.

Malformed knob files are fatal to the ptxas child process. Duplicate assignments follow a last-wins policy: the later command overwrites the earlier runtime value. Identifier matching is case-insensitive from the user's point of view.

tileiras runs no preflight check that the path exists, contains [knobs], or uses valid identifiers. Every knob-file diagnostic comes from ptxas and surfaces through the normal subprocess diagnostic buffer.

Related pages

Driver Overview covers how the produced kernel directives travel into the relocatable object; Driver CLI Options catalogues the user-visible flags that map into pipeline options; ptxas Handoff Protocol documents the ptxas-side knob-file grammar in detail; Attribute System and Lowering documents the full lifecycle of each launch-shape attribute from frontend hint through nvvm.* directive carrier, including which transitions silently drop the fact and produce a degraded kernel.

Subprocess Harness

Abstract

One POSIX subprocess harness drives every external tool tileiras invokes. The same launcher serves ptxas, nvdisasm, and any helper tool — argument construction is tool-specific but process behavior is shared: argv/envp setup, stdio redirection, optional timeout, child status decoding, stderr capture, and resource-usage collection.

The harness itself is CUDA-agnostic. CUDA-specific arguments like --input-as-string, --knobs-file=<path>, --nv-host=<path>, and temporary cubin paths are assembled by the driver before it calls the generic launcher.

POSIX subprocess launcher

The launcher accepts an executable path, argv vector, envp vector, redirection descriptors for stdin/stdout/stderr, optional timeout information, and output buffers for diagnostics and resource usage.

The fork path uses shell-compatible exit codes for exec failure — 127 means the program was not found, 126 means the program existed but could not be executed. The wait decoder maps those codes back into user-facing diagnostics.

int launch_child(ProcessSpec *spec, ChildProcess *child, string *error) {
    if (!spec->setsid && !spec->has_resource_limits)
        return spawn_with_posix_spawn(spec, child, error);

    pid_t pid = fork();
    if (pid < 0)
        return system_error(error, "Couldn't fork");
    if (pid == 0)
        exec_child_or_exit(spec);

    child->pid = pid;
    return 0;
}

stderr-merge optimization

When stderr and stdout target the same sink, the launcher redirects stderr with dup2(stdout, stderr) instead of opening the same destination twice. This is the common ptxas shape — both streams fold into one in-memory diagnostic buffer so the parent can report assembler failure with full context.

SIGALRM wait4 timeout

Timeout handling rides on wait4 and SIGALRM. With a timeout enabled, the parent installs a temporary alarm handler, arms alarm(seconds), and waits on the child. If wait4 is interrupted by the alarm, the parent sends SIGKILL, disarms the alarm, restores the old signal handler, and reaps the child.

int wait_child(pid_t pid, unsigned timeout_s, ProcessResult *result) {
    install_alarm_handler();
    if (timeout_s != 0)
        alarm(timeout_s);

    int status = 0;
    int rc = wait4(pid, &status, 0, &result->usage);
    if (rc < 0 && errno == EINTR && timeout_s != 0) {
        kill(pid, SIGKILL);
        alarm(0);
        restore_alarm_handler();
        waitpid(pid, &status, 0);
        return child_timed_out(result);
    }

    alarm(0);
    restore_alarm_handler();
    return decode_child_status(status, result);
}

Status decoding follows normal POSIX rules. A signaled child reports the signal name and whether a core file was produced. Exit code 126 means the program could not be executed; exit code 127 means command-not-found. Other codes are returned directly to the caller.

ptxas launcher

The ptxas adapter assembles a fixed argv shape around the serialized PTX text:

The normal ptxas call takes the posix_spawn fast path and merges stdout and stderr into one capture buffer. The assembled cubin comes back through the child's stdout, not via a temporary output file.

nvdisasm driver

The SASS dump adapter takes a command string, splits it into argv words, and expects the first word to resolve through PATH. The documented default is:

nvdisasm -c

The adapter writes the ptxas-produced cubin to a temporary file, appends that file path as the final argv element, launches the command, captures stdout, and removes the temporary cubin if the driver created it. When the caller provided an existing cubin path, lifetime management stays with the caller.

int dump_sass(const DumpOptions *opts, ByteSpan cubin, string *out) {
    Argv argv = split_command(opts->dump_sass_command);
    if (argv.len == 0)
        return error("Please provide a valid dumpSassCommand. For example, `nvdisasm -c`.");

    TempFile temp = write_temp_cubin(opts->input_base_name, cubin);
    argv_push(&argv, temp.path);

    int rc = run_child(argv, CAPTURE_STDOUT_AND_STDERR, out);
    remove_temp_file(&temp);
    return rc;
}

Dump-command failures surface as compile-path failures because the driver treats SASS dumping as part of the requested output action.

nvdisasm command-line construction

Once tileiras finishes its MLIR pipeline and PTX emission, the compile dispatcher sub_57A8E0 shells out to two external programs: ptxas (PTX to SASS) and nvdisasm (SASS to disassembled text, embedded as an ELF section). Both invocations route through the subprocess helpers sub_44A36C0 and sub_44A3430, with the raw-ostream sink sub_6CF9C0 draining the child's stdout back into the parent.

Every nvdisasm invocation starts with the literal 5-byte flag block "-uumn" stored at rodata 0x57BB97. The literal packs four flags into one argv token: -u (unsigned offsets), -u (literal second occurrence, triggering the canonical re-entry behaviour known from nvdisasm pre-13.1), -m (mnemonic-only emission), and -n (no header). All five bytes are pushed as one argv element rather than four separate flags.

With config.sanitize == 1 — the only currently-defined sanitize mode, memcheck — the helper appends a 41-byte tail starting with a leading space: -sanitize=memcheck -g-tmem-access-check. The trailing flag -g-tmem-access-check instruments tensor-memory (TMEM) loads and stores, a Blackwell-and-newer concern consistent with the sm_100..sm_121 target table. The full sanitize-on argv therefore consists of the 5-byte "-uumn" token followed by the 41-byte tail token.

The dispatcher writes the literal "nvdisasm -c" at rodata 0x57B6BB as the command-line prefix before the flag block. Program path resolution leaves nvdisasm to PATH; the -c flag asks nvdisasm to emit a section-friendly format suitable for embedding inside the final ELF relocatable.

The sibling ptxas invocation assembles its argv differently. The prefix is "ptxas", followed by --input-as-string and the PTX text as a sized string, then --knobs-file= with the optional knobs path from the env-var registry, and finally --nv-host= with the host triple. The boundary-spec page ptxas Handoff Protocol covers the ptxas side in detail; this section sticks to how the parent assembles the nvdisasm argv.

void buildNvdisasmCmd(const TileirasProgram *p, ArgvBuilder *out) {
    argvPush(out, "nvdisasm");                                          // 0x57B6BB
    argvPush(out, "-c");                                                // 0x57B6C8
    argvPush(out, "-uumn");                                             // 0x57BB97 — 5 bytes, single token
    if (p->sanitize == 1) {
        argvPush(out, " -sanitize=memcheck -g-tmem-access-check");      // 0x57BB9C — leading space included
    }
}

With the argv vector built, sub_44A3430 forks via posix_spawn and wires stdout and stderr through pipes back to the parent. The parent drains the disassembly text from the stdout pipe via sub_6CF9C0 (the raw_ostream sink) and concatenates the captured bytes into the final ELF relocatable as a .nvdisasm section. The temporary cubin path written by the SASS dump adapter is passed as the final argv element, exactly as the previous section described.

Related pages

ptxas Handoff Protocol documents the ptxas side of the boundary including the PTX text protocol and cubin return path; Host Launch ABI and ptxas Knobs covers the --knobs-file= grammar consumed by ptxas; Driver Env Vars and Runtime Gates catalogues the env-var registry that produces the optional PTX_KNOBS_PATH forwarded into the ptxas argv; Driver CLI Options documents the --sanitize=memcheck option that adds the nvdisasm sanitize tail.

TILEIR_CALLBACKS ABI

Abstract

Tileiras emits a callback ABI that lets a runtime shim discover and patch TileIR launch hooks by symbol name. The ABI is not debug logging — it is a compile-time-inserted module table, a pre-load trampoline, a per-function callback table, and an argument-change callback the runtime can invoke when kernel launch arguments or TMA descriptors change.

The public symbol set is:

Global linkage tuples

Every __CUDA_TILEIR_* global is created with an (isConstant, linkage) tuple at definition time. The driver loader inspects those tuples to decide which symbols it may claim and which slots are mandatory. Linkage codes are LLVM's: 0 denotes External (strong), 7 denotes Weak / WeakODR (optional override).

The 64-byte size of __CUDA_TILEIR_FUNC_CALLBACKS is fixed at the type level. The pass emits one such global per kernel function in the module, and the struct shape inside each global is constant. An older note that called the layout variable-size with one slot per emitted func-callback is wrong; the struct shape is pinned and populated by exactly eight insertvalue operations at indices 0..7.

Callback vector ABI

The module-level __CUDA_TILEIR_CALLBACKS global is the entry vector — a constant external object with nine 64-bit slots, total size 72 bytes. The current ABI revision is the integer 0x40, decimal 64. The binary stores this value as the immediate 0x40; the two forms are numerically identical and both are correct references to the same revision.

typedef struct TileirCallbackVector {
    uint64_t init_fn;
    uint64_t fini_fn;
    uint64_t compile_begin_fn;
    uint64_t compile_end_fn;
    uint64_t abi_revision;      /* = 0x40 */
    uint64_t reserved_zero;
    uint64_t mul_multiplier_a;
    uint64_t mul_multiplier_b;
    uint64_t sentinel_zero;
} TileirCallbackVector;

Slots 6 and 7 are MUL multipliers, not flags

Earlier wiki revisions treated slots 6 and 7 as OR-able bit flags. They are not. The body emitter at sub_870430 multiplies a runtime counter by the slot value and writes the product back into the lowered IR. The multiplication routes through sub_868170, the llvm.mul helper. In the generated IR the operation appears as

%v = mul i64 %counter, <slot value>

with <slot value> taken verbatim from slot 6 or slot 7. Treating these slots as OR-combined flag bits would silently miscompile every program that sets more than one — a | b and a * b agree only on single-bit values. Reimplementations must preserve the multiplicative semantics.

Body emission

Three sub-emitters cooperate to produce the callback objects:

• sub_8689C0 (4 293 B) emits the per-revision constant data block — the nine i64 slots for revision 0x40, including the addresses of __CUDA_TILEIR_FUNC_CALLBACKS and __CUDA_TILEIR_ON_PRE_LOAD.
• sub_86DAD0 (~22 KB) emits the full callback dispatch trampoline including the type converter and the nvvm.kernel attribute lift. The related cute.kernel -> nvvm.kernel rename is performed by the downstream D08 pattern CuteKernelToNvvmRewrite at sub_1698C20, documented in TileAS to LLVM Lowering; it is not part of the callback ABI itself.
• sub_870430 is the body emitter that wires the MUL multipliers from slots 6 and 7 into the dispatch path via sub_868170 (the llvm.mul helper).

ON_PRE_LOAD trampoline

At each TileIR launch site, the compiler can emit a pre-load callback call. The runtime-patched symbol is weak and null-guarded, so binaries remain executable even when no runtime shim has installed a callback.

void maybe_call_on_pre_load(void *arg_desc, int64_t sm_num, void *tma_arena) {
    TileirOnPreLoadSlot *slot = &__CUDA_TILEIR_CALLBACKS_ON_PRE_LOAD;
    if (slot == NULL || slot->fn == NULL)
        return;

    slot->fn(arg_desc, sm_num, tma_arena);
}

The argument descriptor is an 11-slot, 88-byte block. The callback receives the descriptor pointer, the sign-extended SM number, and a pointer to the TMA descriptor arena.

Callback signatures

__CUDA_TILEIR_FUNC_ON_ARGUMENTS_CHANGE returns i32, not void. Its first three arguments are pointers — a runtime cookie, an argument-buffer pointer, a TMA-arena pointer. The user kernel arguments follow that prefix in their lowered kernel ABI order.

TMA descriptor shape

Each TMA argument occupies 64 bytes — eight 64-bit words. That shape matches the SM90+ tensor-map descriptor used by cp.async.bulk.tensor. Debug printing splits the descriptor into two four-word rows for readability, but the storage object is one contiguous 64-byte descriptor.

typedef struct TileirTmaDescriptor {
    uint64_t word[8];
} TileirTmaDescriptor;

typedef struct TileirTmaArena {
    uint64_t header[2];
    TileirTmaDescriptor descriptor[8];
} TileirTmaArena;

Device-side descriptor storage must be 64-byte aligned. The compiler relies on the runtime to provide a correctly aligned arena. Non-SM90 targets take the zero path and do not require a populated TMA arena.

Version and backwards-compatibility handling

Revision tracking lives in slot 4 of the module vector. The compiler writes revision 0x40 (= 64 decimal) unconditionally. The runtime checks both the sentinel zero in slot 8 and the revision in slot 4 before installing hooks. No versioned alias symbols like __CUDA_TILEIR_CALLBACKS_v1 exist, so every compatibility check has to go through the vector contents.

The weak runtime-patched pre-load slot and the null guard form the primary backwards-compatibility mechanism. Older runtimes can leave the slot unresolved or null, and the compiled launch site simply skips the callback. The per-function callback table is (1, 0) const+External and fixed at 64 bytes; absent specialized hooks, the argument-change hook is the default callback target.

Related pages

Host Launch ABI and ptxas Knobs covers the kernel-attribute side of launch metadata that the per-function callback table sits alongside; TileAS to LLVM Lowering documents the downstream CuteKernelToNvvmRewrite pass that finalises the nvvm.kernel attribute the callback ABI references; Driver Overview frames where callback emission sits in the larger compile sequence.

MLIR Bytecode Format

Abstract

tileiras consumes a private TileIR bytecode container for cuda_tile modules. The format borrows MLIR bytecode's spirit — unsigned LEB128 integers, section-local offset tables, indexed cross-section references, self-contained attribute payloads — but it is not upstream MLIR bytecode. The magic, version block, section enum, type tags, attribute tags, and cuda_tile opcode space are all private to this binary.

Wire-format divergence vs upstream MLIR. The shipped AttrTag numbering inside sub_59F100 breaks wire compatibility with upstream mlir/Bytecode/BytecodeEnums.h::AttributeTag. Tag 1 is StringAttr here (upstream IntegerAttr=1), tag 4 is DenseElementsAttr (upstream TypeAttr=4), tag 5 is DenseElementsAttr<string> (upstream StringAttr=5), and tag 13 is AssumePredicateAttr — a slot upstream does not define at all. Bytecode emitted by stock MLIR with stock numbering decodes wrong through this binary; bytecode emitted by this binary cannot be consumed by stock MLIR. Any external tool that needs to round-trip through both must freeze the tileiras tag assignments rather than the upstream header. The full 13-tag table sits below in Self-Contained Attribute Dispatch. At the envelope level, the partition signal is magic byte 7 — 0x00 here, '\n' upstream.

Reader contract. The bytecode reader is an in→out transform: a byte buffer goes in, a builtin.module containing a cuda_tile body comes out. Failure paths return nullptr/false with one of the verbatim diagnostics enumerated in the dispatcher sections below.

ModuleOp readTileIRBytecode(ByteSpan input, MLIRContext *ctx) {
    if (!sub_5838A0_validate_header(input))                 // magic + version + dialect-list + blob preamble
        return nullptr;                                     // see "Header Parser (sub_5838A0)" below
    SectionSpans spans = scan_section_table(input);
    StringTable   strings   = read_string_section(spans.string);
    TypeTable     types     = read_type_section(spans.type, ctx);
    ConstantTable constants = read_constant_section(spans.constant, ctx);  // routes payloads through sub_59F100
    DebugTable    debug     = read_debug_section(spans.debug, ctx);        // routes payloads through sub_589B90
    ModuleOp      module    = create_builtin_module(ctx);
    read_globals(spans.global,  module, strings, types, constants);
    read_functions(spans.func,  module, strings, types, constants, debug); // bodies invoke sub_5B13D0 per op
    return module;
}

The rest of the page walks the container at reimplementation depth: file envelope, the six payload sections, the inter-section reference model, validation behavior, and the separate LLVM bitcode path used when tileiras hands an NVVM module to libNVVM.

Overall File Format

The container starts with an 8-byte magic. The first three bytes are the upstream MLIR-bytecode framing prefix (0x06 0x03 0x80); bytes 3–6 spell "Tile"; byte 7 is 0x00 — the tileiras / upstream MLIR split (upstream writes "\nMLIR" starting at byte 7):

06 03 80 54 69 6c 65 00    // MLIR-bc framing + "Tile" + 0x00 terminator

The trailing zero is a sentinel byte inside the magic, not a C-string terminator. Three unsigned-LEB128 VarInts follow — major, minor, optional patch — encoding the Tile version triple. The CUDA 13.1 reader accepts Tile version 13.1.x for any patch value. A mismatched major or minor produces an unsupported Tile version ... diagnostic rather than falling back to upstream MLIR bytecode parsing. See Header Parser (sub_5838A0) for the full magic and version table.

After the version block comes a sequence of sections. Each section starts with one ID byte: the low seven bits are the section ID, the high bit signals an alignment field in the header. A zero ID marks end-of-bytecode.

typedef struct SectionHeader {
    uint8_t id_and_alignment_flag;
    uleb128 length;
    optional<uleb128> alignment;
    uint8_t padding[];
    uint8_t payload[];
} SectionHeader;

length covers the optional alignment field, padding, and payload. Padding bytes are 0xcf. The end marker carries no length, alignment, padding, or payload, and its high bit must be clear.

The reader scans section headers first and records the byte span for each section, then decodes bodies in a second pass. Physical section order is therefore flexible — required sections must exist and all offsets must land inside captured spans, but the order on disk is the producer's choice.

VarInt Encoding

Every multi-byte integer in the container — section length, offset, type reference, operand index, opcode — uses the same unsigned LEB128 variant as upstream MLIR bytecode. The encoding rule is a leading-byte trick: the number of low-order 1-bits in the first byte indicates how many additional bytes follow, the bits above the run-of-ones in that byte form the low bits of the integer, and the subsequent bytes contribute eight bits each in little-endian order.

The leading byte counts its low-order 1-bits, masks them off, and shifts the surviving high bits up to occupy the bottom of the decoded integer; the remaining bytes are appended above those bits. A 10-byte encoding is rejected as overlong — the canonical 9-byte form covers the entire 64-bit range. Signed payloads (the location_index slot in particular) wrap the unsigned VarInt in zig-zag: (n << 1) ^ (n >> 63) going out, (u >> 1) ^ -(u & 1) coming back.

Three concrete decoded examples make the bit pattern unambiguous:

0x0a            → first byte 0000_1010, low bit clear → 1-byte form,
                  payload = 0a >> 1 = 5

0x01 0x02       → first byte 0000_0001, trailing "01" → 2-byte form,
                  payload bits = (byte0 >> 2) | (byte1 << 6)
                              = (0x01 >> 2) | (0x02 << 6)
                              = 0 | 0x80 = 128

0xfb 0xff 0x0f  → first byte 1111_1011, trailing "011" → 3-byte form,
                  payload bits = (byte0 >> 3) | (byte1 << 5) | (byte2 << 13)
                              = (0xfb >> 3) | (0xff << 5) | (0x0f << 13)
                              = 0x1f | 0x1fe0 | 0x1e000 = 0x1ffff = 131071

Producers must always emit the canonical (shortest) encoding for a given integer: an overlong but otherwise-valid encoding decodes to the same integer but flags as "non-canonical VarInt" and rejects the section.

Section Walker Algorithm

Once sub_5838A0 has accepted the envelope and built the blob-section descriptor array, the top-level driver invokes the section walker. The walker is not a recursive parser — it is a fixed, dependency-ordered dispatch over the descriptor array. Earlier sections build the lookup tables that later sections cross-reference, so the order is required even though the on-disk order is the producer's choice.

⚡ QUIRK — on-disk section order is free, walker order is fixed The bytecode envelope lets the producer write sections in any order, but the reader's walk_sections always dispatches them in the fixed dependency-ordered sequence (STRING → TYPE → ATTRIBUTE → IR → optional RESOURCE/DEBUG). Two byte-different files can therefore decode to identical IR, and the reader uses descriptor-array lookup rather than position to find each section. A reimplementation that follows on-disk order to drive parsing will read references against half-built tables and emit unknown <kind> errors for files the official reader accepts.

ParseResult walk_sections(BytecodeReader *r, const BlobSectionDesc *desc, size_t n_desc,
                          MLIRContext *ctx, ModuleOp *out_module) {
    SectionSpans spans = {0};
    for (size_t i = 0; i < n_desc; ++i)
        spans.by_kind[desc[i].section_kind] = (Span){desc[i].offset, desc[i].length};

    /* Sections must be present in dependency order:
     * 1. STRING section (referenced by every later section)
     * 2. DIALECT section (declares the dialects this bytecode uses; implicit here:
     *    the dialect list lives in the envelope, not in its own section)
     * 3. TYPE section (references STRING)
     * 4. ATTRIBUTE / CONSTANT section (references TYPE)
     * 5. IR section: FUNC + GLOBAL records (references all of the above)
     * 6. RESOURCE section (optional)
     * 7. DEBUG section (optional, references the IR section via location slots)
     */
    StringTable    strings   = read_string_section   (r, spans.by_kind[SEC_STRING]);
    TypeTable      types     = read_type_section     (r, spans.by_kind[SEC_TYPE],     ctx, strings);
    ConstantTable  constants = read_constant_section (r, spans.by_kind[SEC_CONSTANT], ctx, strings, types);
    DebugTable     debug     = spans.by_kind[SEC_DEBUG].length
                               ? read_debug_section  (r, spans.by_kind[SEC_DEBUG],    ctx, strings, types)
                               : empty_debug_table();

    ModuleOp module = create_builtin_module(ctx);
    read_globals  (r, spans.by_kind[SEC_GLOBAL], module, strings, types, constants);
    read_functions(r, spans.by_kind[SEC_FUNC],   module, strings, types, constants, debug);

    *out_module = module;
    return success();
}

The walker traverses each section by seeking the reader cursor to span.begin, decoding records until the cursor reaches span.begin + span.length, then asserting that no bytes were left over. A short read inside a section emits a per-section truncation diagnostic; a long read — cursor past span.length after the final record — emits a section-overflow diagnostic. Both fire before any cross-section index from that section is exposed to the next reader, so a corrupt section never poisons downstream lookups with a half-populated table.

The reader trusts the descriptor's offset and length pair to disjoint sections — overlapping spans are not checked. The validation that prevents two sections from claiming the same bytes happens earlier, inside sub_5838A0's preamble loop, where each new descriptor's [offset, offset+length) range is checked against the union of previously accepted ranges.

Header Parser (sub_5838A0)

sub_5838A0 parses the bytecode header. It is invoked from sub_57FF40 once the top-level reader confirms the input buffer starts with the tileiras magic, and it validates the magic prefix, the Tile version field, and the dialect-list / blob-section preamble before handing control to the per-section sub-readers. Everything downstream — section dispatch, attribute decoding, opcode decoding, debug decoding — assumes the header parser has already consumed and accepted this envelope.

Magic prefix. Every TileIR bytecode container opens with an 8-byte magic prefix, compared byte-for-byte against the static literal at rodata 0x45EBF08. The bytes flow through sub_57F420, the bounded-read helper that refuses to walk past the end of the input buffer.

The trailing 0x00 is the byte that separates tileiras-MLIR-bytecode from upstream MLIR-bytecode at the magic level. Upstream MLIR fills the same slot with "\nMLIR", so the shared MLIR-bytecode framing prefix combined with the tileiras-specific terminator cleanly partitions the two dialects. On mismatch the parser emits "invalid magic number at position " concatenated with the buffer offset, then ", got " concatenated with the offending byte sequence, then " expected " concatenated with the expected literal — three verbatim fragments joined into one diagnostic.

Tile version table. Three VarInts follow the magic and encode the Tile version triple. The minimum and maximum supported versions live in a static table at rodata 0x45EBF10:

static const TileVersion supported_versions[] = {
    /*min:*/ { .major = 13, .minor = 1, .patch = 0          }, // inclusive
    /*max:*/ { .major = 13, .minor = 1, .patch = UINT32_MAX }, // inclusive (only 13.1.x)
};

Major and minor VarInts are mandatory. The patch VarInt is optional and defaults to zero when absent — the parser reads it for forward compatibility but never gates on its value. The version-range predicates are:

Only 13.1.x decodes. On reject, the parser emits "unsupported Tile version " followed by the parsed version (formatted via "{0}.{1}.{2}" when the patch field was present, or "{0}.{1}" when absent), then "this reader supports versions [" followed by the rendered min..max range. The two-string format mirrors the patch-optional encoding: a producer that omits the patch VarInt sees its rejection echoed back without a synthetic .0 suffix.

Dialect list and blob preamble. A VarInt dialect-count comes next, then for each dialect a VarInt length-prefixed string followed by a VarInt op-count. Each dialect string is checked against the reader's registered dialect set; unknown dialects emit "unregistered dialect: " concatenated with the offending name. The blob-section preamble that follows is a fixed 24-byte record per section:

typedef struct BlobSectionDesc {
    uint8_t  section_kind;   // one of the seven section IDs documented above
    uint8_t  pad[3];         // alignment padding
    uint64_t offset;         // byte offset from the start of the buffer
    uint64_t length;         // payload length in bytes
    uint32_t alignment;      // required alignment of the payload base
} BlobSectionDesc;

The descriptor array lets the reader build the section-span table without scanning section headers in order: every section's location is known once the preamble is consumed, so the per-section sub-readers can run in any order the driver finds convenient.

Section alignment / ordering invariants. Two structural invariants carry verbatim diagnostics. The end-of-bytecode marker section must have alignment == 0 — otherwise the parser emits "end section should not have alignment flag set". The end marker must also be the last section in the preamble; otherwise the parser emits "end section is not the last section". Both checks fire before any per-section body is decoded, so a malformed envelope is rejected before the reader commits a single byte of section-payload allocation.

Pseudocode shape. The full parser is large, but its skeleton is a magic-and-version prefix followed by the dialect/blob loop. The prefix half looks like:

LogicalResult parseBytecodeHeader(Reader *r, TileVersion *out_ver) {
    uint8_t magic[8];
    if (!readBytes(r, magic, 8))           return emit("failed to read magic"), failure();
    if (memcmp(magic, kMagicLit, 8) != 0)  return emitMagicMismatch(magic),     failure();

    TileVersion v;
    if (!readVarInt(r, &v.major))          return emit("failed to read major"), failure();
    if (!readVarInt(r, &v.minor))          return emit("failed to read minor"), failure();
    if (!readOptionalVarInt(r, &v.patch))  v.patch = 0;
    if (!versionInRange(v))                return emitVersionReject(v),         failure();

    *out_ver = v;
    return success();
}

The dialect-list and blob-preamble loops follow the same shape: every VarInt read pairs with a verbatim diagnostic, every structural invariant is checked before the next field is consumed, and the function returns failure() on the first mismatch so the surrounding driver surfaces a single precise error rather than a cascade.

Per-Section Grammar

The seven section IDs are:

Function records are the richest records in the container:

typedef struct FunctionRecord {
    uleb128 name_string_index;
    uleb128 function_type_index;
    uint8_t flags;
    uleb128 location_index;
    optional<Attribute> optimization_hints;
    uleb128 body_length;
    uint8_t body[body_length];
} FunctionRecord;

The low three flag bits encode visibility, entry/kernel kind, and whether optimization_hints is present. Function bodies contain regions, blocks, block arguments, and operation records. Each operation record starts with an opcode and a location reference, followed by opcode-specific operands, attributes, regions, and result type references.

typedef struct RegionBody {
    uleb128 block_count;
    Block blocks[block_count];
} RegionBody;

typedef struct Block {
    uleb128 arg_count;
    uleb128 arg_type_index[arg_count];
    uleb128 op_count;
    OperationRecord ops[op_count];
} Block;

typedef struct OperationRecord {
    uleb128 opcode;          // cuda_tile opcode, 0..109 in CUDA 13.1
    sleb128 location_index;  // -1 means unknown location
    uint8_t payload[];       // decoded by the opcode-specific reader
} OperationRecord;

Cross-section references are width-sensitive. String and Type offset tables use u32 slots; Constant offsets and some Debug cross-reference tables widen to u64 so large dense payloads remain addressable. Function and Global records are sequential rather than table-indexed.

VarInts are unsigned LEB128 capped at 10 bytes — overlong 11-byte encodings are rejected. Signed integer payloads use zig-zag decoding; APInt multiword payloads apply zig-zag per word.

Type Tag Dispatch

sub_59C710 decodes the Type-section payload table. The routine is 6 474 bytes long and switches on a one-byte TypeTag at 0x59C79D; it is the type-side sibling of the Op, Attr, and Debug dispatchers covered later. Each type-table entry triggers one call from sub_58C0C0 (the cached Type-by-index lookup), itself reached via sub_58C400 whenever a downstream consumer needs to resolve a type reference. The TypeTag numbering is independent of upstream MLIR's BytecodeTypeOpcodes.td and uses the dense 0..18 assignment shown below.

Type records start with a TypeTag followed by a tag-specific payload:

Tile types carry an element type and shape. Tensor views add strides. Partition views point at a tensor-view type, then add tile shape, dimension map, and partition mode. Function types carry parameter and result type-index vectors. The per-tag operand layout in wire order is:

Type read_type(Reader *r) {                                  // sub_59C710
    uint64_t tag = read_uleb128(r);                           // (1) one TypeTag VarInt

    switch (tag) {
    case TYPE_I1: case TYPE_I8: case TYPE_I16: case TYPE_I32: case TYPE_I64:
        return integer_type(width_for_tag(tag));
    case TYPE_F16: case TYPE_BF16: case TYPE_F32: case TYPE_TF32:
    case TYPE_F64: case TYPE_F8E4M3FN: case TYPE_F8E5M2:
        return float_type_for_tag(tag);
    case TYPE_POINTER:
        return pointer_type(read_type_ref(r));
    case TYPE_TILE:
        return tile_type(read_type_ref(r), read_i64_shape(r));
    case TYPE_TENSOR_VIEW:
        return tensor_view_type(read_type_ref(r), read_i64_shape(r), read_i64_strides(r));
    case TYPE_PARTITION_VIEW:
        return partition_view_type(read_type_ref(r), read_shape(r), read_dim_map(r), read_partition_mode(r));
    case TYPE_FUNCTION:
        return function_type(read_type_ref_list(r), read_type_ref_list(r));
    case TYPE_TOKEN:
        return token_type();
    default:
        return registered_extension_type(tag, r);
    }
}

Operation Opcode Dispatch

One master dispatcher decodes the entire cuda_tile public opcode space: sub_5B13D0, 10 650 bytes, jump table at 0x5B158E. It takes one operation record at a time, reads the opcode as a VarInt, attaches a source location, switches on the opcode integer, and either inlines the canonical op-builder skeleton or tail-calls a dedicated per-op parser. It returns true if the opcode was recognized and the resulting MLIR Operation* passed verification, false otherwise.

Calls come from sub_57FF40, the bytecode-parse-into-scratch path that walks a function body and stages each operation into per-block operand and location tables before final placement into the materialized region. sub_5B13D0 itself stays op-local: no block-structure work, no region allocation, and no cursor advances outside its argument context.

The canonical body is a five-step sequence repeated for every operation record:

bool parse_operation(BytecodeReader *r, OpBuilder *b, ValueList *operands,
                     LocAttrList *locs, AttrList *attrs, Operation **out) {
    uint64_t opcode;
    if (!sub_5847D0_read_opcode(r, &opcode))                       // (1) read opcode VarInt
        return error(r, "failed to read operation opcode.");

    Location loc;
    if (!read_location(r, locs, &loc))                             // (2) resolve source location
        return error(r, "failed to read operation location.");

    switch (opcode) {                                              // (3) dispatch on opcode 0..109
    case 0x00:  return sub_58C5C0(r, b, operands, attrs, loc);     // (4a) delegate to per-op parser
    case 0x09:  return build_inline_bitcast(r, b, operands, attrs, loc, out);  // (4b) inline skeleton
    /* ... 108 more cases ... */
    default:    return error(r, "unknown or unimplemented opcode: ", opcode);
    }
    /* (5) common LABEL_280 cleanup chain frees scratch buffers and returns the verifier result */
}

Step (1) calls the opcode reader, which returns the VarInt and a position-advanced cursor. Step (2) resolves the location either by reading a LocAttr index from the current location-table slot or, when no location was emitted, by synthesizing an UnknownLoc from the MLIR context. Step (3) is the 110-entry jump table at 0x5B158E. Step (4) splits two ways: a tail call to the per-op parser, which reads its own result types, operand indices, and attributes before calling sub_5847D0; or an inline ODS-shaped skeleton that reads the result type via sub_58C400, reads zero or more operand indices via sub_585AD0, then calls sub_5847D0 with the verbatim mnemonic literal embedded in the dispatcher binary. Step (5) is shared by every case path that does not tail-call: it frees the SSO op-name scratch and the transient attribute vector, then returns the verifier's verdict on the freshly built operation.

Four verbatim diagnostic strings are emitted by this dispatcher and its helpers:

• "failed to read operation opcode."
• "failed to read operation location."
• "unknown or unimplemented opcode: "
• "failed to create operation '…' due to verification error."

The first three live inside sub_5B13D0 itself. The fourth comes from sub_5847D0, the create-and-verify helper every case path eventually reaches; its prefix and suffix wrap the mnemonic the current case passed in, so a verifier rejection on, say, cuda_tile.addf surfaces as failed to create operation 'cuda_tile.addf' due to verification error..

Two reserved opcode ranges sit in the jump table and fall through to the default arm. Opcodes 25–36 inclusive (0x19–0x24) and opcodes 52–57 inclusive (0x34–0x39) are unassigned in CUDA 13.1 and emit the "unknown or unimplemented opcode: " diagnostic if a producer encodes them. These gaps line up with corresponding holes in the public ODS opcode assignment — room for future op additions without invalidating already-shipped bytecode.

The full 110-row dispatch table follows. Each row gives the decimal opcode, the cuda_tile mnemonic, and either the dedicated parser address or inline for cases handled by the inline ODS-shaped skeleton inside sub_5B13D0. Region-bearing ops (entry, for, if, loop, module, reduce, scan) delegate to dedicated parsers because they additionally consume nested region bodies before calling sub_5847D0.

Opcode 0x6E — atan2 in the public 13.2 opcode space — is absent from this binary. The dispatcher has no case for it and embeds no cuda_tile.atan2 mnemonic; encoding the op would land on the default arm and surface the "unknown or unimplemented opcode: " diagnostic. Consistent with a 13.1-vintage reader that predates the atan2 addition.

Worked encode example. Take the operation

%c = cuda_tile.addi %a, %b : tile<8 × i32>

and assume the surrounding context has already populated the per-section tables so that %a is value-table entry 4, %b is value-table entry 5, and tile<8xi32> is type-table entry 3. The function body's operation-record encoder writes seven fields in fixed order, each as a single VarInt:

The final on-wire byte stream for this operation record is therefore exactly seven bytes:

03 7f 03 02 04 05 00

A run with --lineinfo replaces the 0x7f sentinel with a non-negative LocAttr index encoded as a positive zig-zagged VarInt — typically one byte (0x00 for index 0, 0x02 for index 1, 0x04 for index 2, and so on) — and stretches the record to eight bytes. A run with a non-empty inline attribute dictionary stretches the trailing 0x00 into a VarInt index into the attribute table, again typically one byte for small modules.

The operation cost in the IR section is therefore constant in the number of operands plus a tiny constant for the bookkeeping fields, and is independent of the mnemonic string. The mnemonic cuda_tile.addi lives once — in the dispatcher's per-opcode string literal at dispatch case 0x03 — and never appears in the per-operation byte stream.

The three functions around this dispatcher fit together cleanly. sub_5847D0 is the opcode-reader producing the integer the master switch keys on, and every case of sub_5B13D0 either inlines the ODS skeleton that ends in a sub_5847D0 call or tail-calls a per-op parser that itself ends in sub_5847D0. The Location decoder runs once per operation — between the opcode read and the switch — and writes the resolved Location into the per-op slot every case path reads when populating its OperationState. One layer up, sub_57FF40 is the bytecode-parse-into-scratch path the driver invokes per function body; it calls sub_5B13D0 in a loop for each operation record while maintaining the operand, location, and attribute vectors the dispatcher consumes through its argument context.

Self-Contained Attribute Dispatch

Every attribute payload is self-contained. Constants, function optimization hints, and the inline attribute slots on operations all funnel through the same decoder. The 13-case dispatcher inside sub_59F100 recognizes string, float, type, dense-elements (int/float and string variants), divisibility, dense-i64-array (two layout variants), same-elements, bounded (three discriminator variants), and assume-predicate attributes. The integer, bool, array, dictionary, and optimization-hints attribute kinds are not handled by this dispatcher — they arrive through the upstream MLIR builtin dispatcher path on a different code path. A Global initializer must resolve to a dense integer-or-floating elements attribute even though the Constant section can store a broader attribute set.

Anywhere the reader encounters an attribute payload that does not come pre-resolved through the Constant offset table — operation attribute dictionaries, type-attribute slots, location-attribute slots, the Constant payloads themselves — the bytes route through sub_59F100. This dispatcher is the attribute-side sibling of the 110-case opcode switch. Roughly 8 KB, it dispatches on a uint32_t AttrTag through a jump table at the entry switch and returns either a heap-allocated Attribute on success or nullptr on failure. The caller pushes the result into the bytecode reader's attribute table; failures propagate up to the section-level error path that aborts the load.

The shipped tileiras tag numbering is wire-format-breaking versus upstream MLIR. The two numberings are reproduced side by side so the divergence is unambiguous:

The only tag that matches upstream by accident is tag 2 (FloatAttr in both). Every other tag in the 1..13 range disagrees: tag 1 is StringAttr here versus upstream IntegerAttr, tag 4 lands on DenseElementsAttr instead of TypeAttr, tag 5 lands on DenseElementsAttr<string> instead of StringAttr, tag 6 lands on DivByAttr instead of ArrayAttr, and so on. Going the other direction, an AssumePredicateAttr emitted by tileiras at tag 13 has no destination in upstream's table at all — upstream's reader rejects the tag with its own default-arm diagnostic.

The structural consequence is sharper than tag-by-tag remapping: tileiras's bytecode reader cannot consume upstream MLIR's bytecode files when those files carry attributes, and tileiras-emitted bytecode (when a future build links a writer) cannot be loaded by stock MLIR. The textual MLIR asm is still interoperable through the printer / parser, but the bytecode wire format is a hard fork. Any external tool that wants to round-trip MLIR bytecode through both tileiras and upstream MLIR must freeze the tag assignments used by this binary rather than the ones in the upstream header. The upstream numbering is reserved for future stock cuda_tile builds; the shipped binary stays compatible with an earlier frozen scheme.

The thirteen recognized tag values, the attribute kinds they construct, and the per-tag builder functions are:

Tags 1, 2, and 3 decode inline in sub_59F100 itself. Tag 1 resolves a string via sub_59AD90 and wraps it in StringAttr. Tag 2 reads a type reference, validates it as a FloatType via sub_58C400, reads an inline APFloat payload via sub_586200, and dispatches through the sub_4462700-family float-type-builder to produce a FloatAttr. Tag 3 reads a type reference via sub_58BDE0 and wraps it in TypeAttr. Every other tag tail-calls a dedicated sub-decoder in the sub_59FB80–sub_5A0620 cluster; those decoders read the tag-specific payload, build the corresponding attribute, and either return the new attribute or emit the per-decoder error string and return nullptr. The default arm covers tag 0 and every tag above 13: it emits "unsupported AttributeTag " (verbatim, trailing space included) concatenated with the tag integer and the suffix " for self-contained attribute", then returns nullptr. Diagnostics route through the standard emitter chain sub_57EA50 / sub_581460 at severity 0x103.

The canonical body is the entry prologue plus the 13-way switch and the default-arm error path:

Attribute *parseSelfContainedAttribute(BytecodeReader *r) {
    uint32_t tag;
    if (!read_varint_u32(r, &tag)) {                              // (1) read AttrTag VarInt
        emit_error(r, "failed to read AttributeTag for self-contained attribute.");
        return NULL;
    }

    switch (tag) {                                                // (2) 13-case dispatch
    case  1: return read_string_attr_inline(r);                   // StringAttr via sub_59AD90
    case  2: return read_float_attr_inline(r);                    // FloatAttr (type-ref + APFloat)
    case  3: return read_type_attr_inline(r);                     // TypeAttr (type-ref only)
    case  4: return sub_59FB80(r);                                // DenseElementsAttr int/float
    case  5: return sub_59FCD0(r);                                // DenseElementsAttr string
    case  6: return sub_59FE40(r);                                // DivByAttr
    case  7: return sub_59FF60(r);                                // DenseI64ArrayAttr variant A
    case  8: return sub_5A0080(r);                                // DenseI64ArrayAttr variant B
    case  9: return sub_5A01A0(r);                                // SameElementsAttr
    case 10: return sub_5A02C0(r);                                // BoundedAttr variant 0
    case 11: return sub_5A03E0(r);                                // BoundedAttr variant 1
    case 12: return sub_5A0500(r);                                // BoundedAttr variant 2
    case 13: return sub_5A0620(r);                                // AssumePredicateAttr
    default:                                                      // (3) default-arm error path
        emit_error(r, "unsupported AttributeTag ", tag, " for self-contained attribute");
        return NULL;
    }
}

Twenty-one verbatim diagnostic strings are reachable from sub_59F100 and its inline tag arms. They are reproduced below exactly as they appear in the binary; the trailing period or trailing space is part of the string. The "Tag" column identifies which switch arm emits each string, and the "Trigger" column states the failure condition that surfaces the diagnostic.

The "unsupported AttributeTag " / " for self-contained attribute" pair is the canonical sentinel for forward-incompatible bytecode: any future tileiras that adds AttrTag values 14+ will be rejected by this CUDA 13.1 reader with that exact pair of fragments wrapping the offending integer. Producers that need to stay compatible with the shipped binary must restrict themselves to the thirteen tags above.

This dispatcher relates to the rest of the bytecode reader the same way the opcode dispatcher does. Callers are sub_5A0A50 (the Constant-section attribute-table populator), sub_5A1CA0 (the cuda_tile.assume parser, which carries a self-contained AssumePredicateAttr payload), sub_5A2410, and sub_5A7AD0. Each one passes a reader cursor and receives a heap-allocated Attribute or a propagated failure back; no caller attempts to recover from a nullptr return. The per-tag builders in the sub_59FB80–sub_5A0620 cluster share callees with the inline arms — sub_57BCF0 for VarInts, sub_58C400 for Type-by-index, sub_58BDE0 for Attr-by-index, sub_456A580 for vector reservation, sub_586200 for inline APInt/APFloat decoding — so the dispatcher's behavior is fully described by the tag table above plus the diagnostic table.

This section is the attribute-side companion to the Operation Opcode Dispatch above. The corresponding dialect-level question — which cuda_tile attributes are recognized by this binary at all — is summarized in Dialect Bytecode Reader/Writer Status — Status Matrix. The cuda_tile bytecode reference details the attribute-readers per dialect, including the exact payload layout each sub_59FB80–sub_5A0620 builder consumes.

Debug-Info Attribute Dispatch

Debug information — DICompileUnit, DIFile, DILexicalBlock, DILoc, DISubprogram, CallSite — does not flow through the AttrTag dispatcher above. It has its own third dispatcher: sub_589B90. The function is 8 779 bytes, switches on a uint32_t DebugTag at 0x589E05, and fires whenever a DistinctAttr-class or LLVM DI* attribute appears in any Location slot, in the Debug section's parallel attribute table, or — through recursion — as a nested scope reference inside another debug attribute. It rounds out the trio of bytecode dispatchers alongside the 13-case AttrTag switch in sub_59F100 and the 110-case Op opcode switch in sub_5B13D0. All three are reached from the top-level bytecode-parse routine sub_57FF40, which walks function bodies and routes each encountered tag through the appropriate dispatcher based on its section context.

The seven recognized tag values, the attribute kinds they construct, and the per-tag builder strategy are:

Tag 5 is the only case that tail-calls a dedicated sub-parser. DISubprogramAttr carries the heaviest payload of the seven — compile-unit reference, function name string, linkage-name string, enclosing scope reference, line number, generic flags, special-purpose discriminator, and an optional spFlags word — and the body is large enough that the dispatcher delegates rather than inlining it. The 8 779-byte main body of sub_589B90 open-codes every smaller case: tags 1, 2, 3, 4, and 6 each have their full read sequence inline in the switch arm, with one diagnostic per failable field read.

Tags 3, 4, and 6 decode scope-like cross-references and call sub_589B90 recursively. A DILexicalBlock references its enclosing scope, itself another DI* attribute. A DILoc references both its containing scope and, if inlined, an inlined-at location, both of which route back through the same dispatcher. A CallSite references its caller subprogram and its callee subprogram, again as full debug attributes. The recursion has no cycle detection: the bytecode writer never emits cycles and the reader trusts the input, so a malformed stream with a transitive scope cycle recurses until the stack is exhausted. Producers that hand-craft debug bytecode must topologically order the debug table so every reference points at an attribute emitted earlier in the stream.

Worked example — recursive scope cycle. A small input that exercises the recursion (and shows the cycle failure mode) starts from a DILexicalBlock whose scope reference points back at itself. Upstream MLIR forbids self-cycles, but a hand-crafted bytecode stream can emit them; the reader's behavior on such input is observable and worth documenting.

debug attribute table, index 0:
    tag  = 3                          // DILexicalBlock
    scope_ref  = attr_index(0)        // self-reference (the cycle)
    file_ref   = attr_index(1)        // forward ref to the DIFile below
    line       = 42
    column     = 0
debug attribute table, index 1:
    tag  = 2                          // DIFile
    name_string_index      = 7
    directory_string_index = 8

Loading this table walks sub_589B90 like so:

1. Top-level call: tag VarInt 0x03 → DILexicalBlock arm.
2. Read scope-ref VarInt 0x00 → recurse into sub_589B90 at attribute index 0.
3. Top-level frame is still in flight; the recursive call reads tag VarInt 0x03 again and recurses again.
4. The cycle has no fixed point: step (3) repeats until the C stack overflows. Empirically this fires after a few thousand frames on a default 8 MiB stack; the process dies with SIGSEGV, not with a tileiras diagnostic.

A well-formed equivalent of the same intent emits the inner attribute first and indexes into it from the outer one:

debug attribute table, index 0:
    tag  = 2                          // DIFile, emitted first
    name_string_index      = 7
    directory_string_index = 8
debug attribute table, index 1:
    tag  = 3                          // DILexicalBlock, now references a *prior* attr
    scope_ref  = attr_index(0)        // resolves cleanly
    file_ref   = attr_index(0)
    line       = 42
    column     = 0

The constructed mlir::LocationAttr is a DILexicalBlockAttr whose scope and file both point at the DIFileAttr at index 0. The recursion bottoms out on the first call because tag 2 (DIFile) has no scope-shaped fields and decodes inline.

The takeaway is asymmetric: well-formed input always terminates in a single bounded recursion sweep because attributes are emitted in topological order; ill-formed input that introduces a cycle is detected only by stack exhaustion. A future reader that wants to harden this path would maintain a visited-set keyed by attribute index during the recursive walk and emit a "cyclic debug attribute reference at index " diagnostic on revisit. The shipped CUDA 13.1 reader does not.

The canonical dispatcher body is the entry-prologue VarInt read plus the 7-arm switch:

Attribute *parseDebugAttribute(BytecodeReader *r) {
    uint32_t tag;
    if (!read_varint_u32(r, &tag)) {                                  // (1) read DebugTag VarInt
        emit_error(r, "failed to read kind tag");
        return NULL;
    }

    switch (tag) {                                                    // (2) 7-arm dispatch
    case 0:  emit_error(r, "unknown debug attribute tag");            // default-equivalent inline arm
             return NULL;
    case 1:  return read_di_compile_unit_inline(r);                   // DICompileUnitAttr
    case 2:  return read_di_file_inline(r);                           // DIFileAttr
    case 3:  return read_di_lexical_block_inline(r);                  // recurses on scope
    case 4:  return read_di_loc_inline(r);                            // recurses on scope + inlined-at
    case 5:  return sub_588E60(r);                                    // DISubprogramAttr (delegated)
    case 6:  return read_call_site_inline(r);                         // recurses on caller and callee
    default: emit_error(r, "unsupported DebugTag ", tag);             // forward-incompatibility sentinel
             return NULL;
    }
}

Fourteen verbatim diagnostic strings are reachable from sub_589B90 and its inline tag arms. They are reproduced below exactly as they appear in the binary; the trailing punctuation, if any, is part of the string. The "Tag" column identifies which switch arm emits each string, and the "Trigger" column states the failure condition that surfaces the diagnostic.

Two structural observations follow from the table. The tag 4 arm emits its inner errors under the FileLineColLoc name rather than DILocAttr, which reflects how DILocAttr is built on top of MLIR's FileLineColLoc primitive: line and column flow through a sub-helper shared with plain locations, and that sub-helper emits its own diagnostics under its own attribute name. The tag 6 arm spells CallSiteLoc rather than CallSiteAttr for the same reason: the location form predates the attribute form in MLIR's debug-info subsystem, and the embedded diagnostic literals carry the older name. Consumers parsing tileiras error output must accept both spellings as referring to the same sub_589B90 switch arms.

sub_589B90's relationship with the rest of the bytecode reader mirrors the other two dispatchers. Callers are sub_588E60 itself (when the DISubprogramAttr body needs to recurse for its CU or scope reference), sub_58BDE0 (the Attr-by-index cached lookup, which routes any DistinctAttr-class index through the debug dispatcher), and itself recursively as described above. Every caller passes a reader cursor and receives a heap-allocated Attribute or a propagated nullptr; nobody recovers from a failed debug read, so a single corrupt DebugTag aborts the entire module load. Shared callees with the other two dispatchers are sub_57BCF0 for VarInts, sub_58BDE0 for Attr-by-index, the sub_57EA50 / sub_581460 emitter pair at severity 0x103 for diagnostics, and the string-table lookup path that emits the shared "string index " prefix on out-of-range references.

This section is the debug-info companion to the Operation Opcode Dispatch and Self-Contained Attribute Dispatch sections above. The dialect-level question of which debug attributes the bytecode writer actually emits in CUDA 13.1 is tracked in Dialect Bytecode Reader/Writer Status — Status Matrix; the seven tags above correspond to the strict subset of LLVM DI* attributes that survive round-tripping through this reader.

Ordering and Diagnostics

Physical section order stays flexible because the reader captures section spans before decoding bodies. Required sections are String and Func. Structural errors live in a separate channel from mandatory-section errors and per-section decode errors — the distinction lets tools tell "not a TileIR file" apart from "valid TileIR envelope with a malformed Type section".

The driver also distinguishes TileIR bytecode from upstream MLIR bytecode. When the input looks like ordinary MLIR bytecode, the diagnostic says so rather than emitting only a generic magic-number failure.

NVPTX LLVM Bitcode Path

When tileiras takes the in-process libNVVM path, it also serializes LLVM bitcode — a different format entirely from the TileIR bytecode described above. The NVPTX64 data layout is stamped onto the LLVM module unconditionally before serialization:

const char *NVPTX64_DATA_LAYOUT =
    "e-p:64:64:64-p3:32:32:32-i1:8:8-i8:8:8-i16:16:16-i32:32:32-"
    "i64:64:64-i128:128:128-f32:32:32-f64:64:64-v16:16:16-v32:32:32-"
    "v64:64:64-v128:128:128-n16:32:64";

The override discards whatever data layout MLIR's LLVM dialect translation left behind. The NVPTX ABI fact worth pinning down is address space 3 — 32-bit shared memory.

NVPTX modules ship as bare LLVM bitcode, not the wrapper-framed object format used for some host triples. The bitcode identifies itself as LLVM 21-era output and reaches libNVVM under the module name mlir-input.

BitcodeBuffer serialize_for_libnvvm(LLVMModule *m) {
    module_set_data_layout(m, NVPTX64_DATA_LAYOUT);

    BitcodeBuffer bc = write_llvm_bitcode(
        m,
        /*wrapper=*/false,
        /*producer=*/"LLVM21.0.0git");

    return bc;
}

NvvmResult compile_with_libnvvm(BitcodeBuffer bc) {
    NvvmProgram prog = nvvmCreateProgram();
    nvvmAddModuleToProgram(prog, bc.data, bc.size, "mlir-input");
    nvvmCompileProgram(prog, options);
    nvvmVerifyProgram(prog, options);
    return nvvmGetCompiledResult(prog);
}

The default command-line path still goes through PTX and ptxas. The bitcode path only matters when the pipeline is wired to use libNVVM directly. The NVPTX target initialization and the data-layout stamping documented above are covered end-to-end in NVPTX Bring-up and Target Init; the libdevice bitcode that gets linked into the same module is documented in libdevice Overview — Link, inline, simplify.

Cross-References

This page documents the wire-format the bytecode reader consumes; four companion pages cover the reader from complementary angles. Frontend Contract and Tile IR Emission documents the producer-side rules a frontend must satisfy to emit conformant bytecode — the dialect list, the magic and version constants, the AttrTag numbering, and the canonical VarInt encoding — and catalogues the common emission mistakes that produce buffers this reader rejects. Dialect Bytecode Reader/Writer Status restricts the wire format to the dialects that actually ship a reader — cuda_tile is the only TileIR dialect with one, and no TileIR dialect ships a writer — and frames the asymmetry as a deliberate input-only driver contract. Dialect Asm-Printer Status documents the textual side of the same contract, because round-trip workflows on intermediate dialects rely on the asm-printer rather than the bytecode writer this binary does not link. cuda_tile Bytecode Reader zooms back in on the cuda_tile-private dispatchers — the 18-case TypeTag dispatcher, the cuda_tile-specific AttrTag payload shapes that route through the 13-case dispatcher documented above, and the 110-case Op opcode dispatcher whose dispatch table is reproduced in Operation Opcode Dispatch. The wire-format-breaking AttrTag divergence is the most consequential single fact across all four pages: a bytecode file containing attributes is not portable between tileiras and stock MLIR, and any reimplementation must freeze the tileiras numbering reproduced in Self-Contained Attribute Dispatch.

Dialect Bytecode Reader/Writer Status

Abstract

tileiras consumes TileIR bytecode in one direction only. It accepts serialized cuda_tile modules at the driver boundary, lowers them through several internal dialects, and emits PTX or object code — none of those dialects ever round-trip back to MLIR bytecode. The compatibility rule is simple: cuda_tile is the only TileIR dialect with a linked bytecode reader, and no TileIR dialect in this binary ships a bytecode writer. Downstream dialects — nv_tileaa, nv_tileas, cute, cute_nvgpu, cutlass — are in-memory pipeline representations.

Reader architecture

A single reader-driver walks the bytecode container in a fixed order: file envelope (magic, version, dialect-list), string section, type section, attribute section, IR section (functions and globals), resource section, and an optional debug section. The driver reads each section header, validates the recorded byte count against the available span, and hands the section body to a per-section dispatcher. Section dispatchers iterate the body, reading one record at a time and routing each record through a tag-keyed switch onto a typed handler.

The reader is decoder-only at every level. Four wire-format dispatchers carry the work — one for operation opcodes (110 cases over the cuda_tile public opcode space), one for self-contained attribute payloads (13 cases, wire-format-breaking versus upstream MLIR), one for type tags (18 cases), and one for debug attributes (7 cases). None has a sibling writer dispatcher linked in. A reimplementation that wants to produce TileIR bytecode must build its own encoder against the tag numberings documented in MLIR Bytecode Format; the shipped reader is the only source of truth for the wire-format constants, and the attribute-tag numbering deliberately diverges from upstream MLIR (tag 1 is StringAttr, tag 13 is AssumePredicateAttr, magic byte 7 is 0x00).

Status Matrix

Upstream MLIR builtin bytecode support is still linked in because the file container uses MLIR infrastructure for built-in types and attributes. That does not mean the TileIR dialects themselves provide general MLIR bytecode round-tripping.

Reader contract

The cuda_tile reader is the only path that materializes IR at the driver boundary. Its top-level loop validates the envelope, scans the section table, then dispatches each section in container order — the later sections reference indices into the earlier ones, so reordering would break cross-section lookups.

ModuleOp read_tileir_module(ByteSpan input, MLIRContext *ctx) {
    BytecodeReader r = bytecode_reader_init(input);

    /* envelope: 8-byte magic, LEB128 version, dialect-list, blob preamble */
    if (!read_and_verify_magic(&r))         { diag(r, "invalid TileIR magic");   return nullptr; }
    uint64_t version = read_leb128(&r);
    if (!version_is_supported(version))     { diag(r, "unsupported version");    return nullptr; }
    DialectList dialects = read_dialect_list(&r);  /* requires "cuda_tile" entry */

    /* sections appear in a fixed order; later sections index into earlier ones */
    SectionTable sections = scan_section_table(&r);

    StringTable    strings    = read_string_section   (&r, sections.string);
    TypeTable      types      = read_type_section     (&r, sections.type,      ctx, strings);
    AttributeTable attrs      = read_attribute_section(&r, sections.attribute, ctx, strings, types);
    ResourceTable  resources  = read_resource_section (&r, sections.resource,  ctx, strings);
    DebugTable     debug      = sections.debug.present
                                ? read_debug_section  (&r, sections.debug,     ctx, strings, attrs)
                                : empty_debug_table();

    ModuleOp module = create_builtin_module(ctx);
    read_ir_section(&r, sections.ir, module, strings, types, attrs, resources, debug);
    return module;
}

Each section dispatcher follows the same shape — read a small header (record count, dialect index, optional flags), then iterate record bodies, routing each record's lead byte through the appropriate tag switch:

ParseResult read_attribute_section(BytecodeReader *r, SectionSpan span, ...) {
    bytecode_reader_seek(r, span.begin);
    uint64_t count = read_leb128(r);
    for (uint64_t i = 0; i < count; ++i) {
        uint8_t tag = read_byte(r);
        switch (tag) {
            case ATTR_STRING:            parse_string_attr(r, /*has_dialect=*/false); break;  // tag 1
            case ATTR_FLOAT:             parse_float_attr(r); break;                          // tag 2
            case ATTR_TYPE:              parse_type_attr(r); break;                           // tag 3
            case ATTR_DENSE_ELT:         parse_dense_elements_attr(r); break;                 // tag 4
            case ATTR_DENSE_ELT_STRING:  parse_dense_elements_string_attr(r); break;          // tag 5
            case ATTR_DIV_BY:            parse_div_by_attr(r); break;                         // tag 6
            case ATTR_DENSE_I64_ARRAY_A: parse_dense_i64_array_attr_a(r); break;              // tag 7
            case ATTR_DENSE_I64_ARRAY_B: parse_dense_i64_array_attr_b(r); break;              // tag 8
            case ATTR_SAME_ELEMENTS:     parse_same_elements_attr(r); break;                  // tag 9
            case ATTR_BOUNDED_LO:        parse_bounded_attr(r, /*variant=*/0); break;         // tag 10
            case ATTR_BOUNDED_HI:        parse_bounded_attr(r, /*variant=*/1); break;         // tag 11
            case ATTR_BOUNDED_LO_HI:     parse_bounded_attr(r, /*variant=*/2); break;         // tag 12
            case ATTR_ASSUME_PREDICATE:  parse_assume_predicate_attr(r); break;               // tag 13
            default:                     diag_unsupported_attr_tag(r, tag); break;
        }
    }
    return success();
}

The tag-to-attribute-kind mapping in the case names above is the wire-format-breaking tileiras numbering documented in MLIR Bytecode Format — Self-Contained Attribute Dispatch; tags IntegerAttr/BoolAttr/ArrayAttr/DictionaryAttr/OptimizationHintsAttr from upstream MLIR are not present in this dispatcher and arrive instead through the upstream MLIR builtin reader path (builtin dialect's own bytecode arms).

Operation records inside the IR section follow the same dispatch shape with the 110-case opcode switch in place of the 13-case attribute switch. Type records use the 18-case type switch; debug records use the 7-case debug switch. The four switches are independent — they share no fallthrough — and each one terminates a section: when the section's byte count is exhausted, the reader returns to the driver loop.

Non-cuda_tile TileIR bytecode is rejected at the driver boundary. When the input looks like ordinary upstream MLIR bytecode (different magic byte 7, different attribute-tag numbering), the driver reports that shape explicitly instead of silently reinterpreting it as TileIR.

Reader-only contract

The missing writers are user-visible. A tool can hand tileiras a cuda_tile bytecode module and ask for compiled output, but it cannot ask this binary to emit optimized TileIR bytecode or any intermediate-dialect bytecode. The capability surface is a one-line predicate:

bool tileiras_can_read_bytecode (const char *dialect) { return strcmp(dialect, "cuda_tile") == 0; }
bool tileiras_can_write_bytecode(const char *dialect) { (void)dialect; return false; }

The asymmetry is deliberate. Round-trip workflows use textual IR dumps for inspection; cacheable intermediate artifacts require an external writer linked against compatible dialect implementations. Treating the intermediate dialects as encoder-absent on purpose lets the pipeline evolve without freezing a stable wire format for every internal representation — the only stable surface is the cuda_tile input boundary.

Several driver behaviors fall out of this asymmetry: the command-line input must be TileIR bytecode (not generic MLIR bytecode); the driver exposes no --emit-bytecode or --write-bytecode mode; intermediate IR dumps, when enabled, are textual MLIR asm rather than bytecode; the internal dialect stack can change shape between releases without breaking external tooling.

Cross-references

The detailed wire format consumed by the reader-driver — file envelope, section ordering, every tag enumeration, the validation diagnostics, and the LLVM bitcode path used for NVVM modules — lives in MLIR Bytecode Format. The textual-asm side of the reader contract, including how intermediate dialects are inspected when bytecode serialization is unavailable, is covered in Dialect Asm-Printer Status. The dialect-level semantics that the bytecode reader materializes are documented in the per-dialect references — cuda_tile bytecode reference and cuda_tile Overview for the input dialect, plus the corresponding overviews for nv_tileaa, nv_tileas, cute, cute_nvgpu, and cutlass.

Dialect Asm-Printer Status

Abstract

An MLIR asm printer turns in-memory operations, attributes, types, and values into the textual .mlir form used for human inspection and crash dumps. Each dialect contributes to the result through three mechanisms: a printType / parseType pair that handles dialect-specific type bodies, a printAttribute / parseAttribute pair that handles dialect-specific attribute bodies, and an OpAsmDialectInterface that supplies short readable aliases plus per-value SSA-name hints. Per-operation pretty-printing is layered on top through OpAsmOpInterface hooks attached at ODS time. When a dialect leaves a hook unimplemented, MLIR's default trampoline takes over and emits the verbose generic form — "dialect.op"(operands) : (types) -> types for operations, !ns<...stored_payload...> for types, #ns<...stored_payload...> for attributes.

Textual assembly is the only inspection path for the non-input dialects inside tileiras, since this binary never serializes them as bytecode. The printer surface is intentionally uneven: dialects near the user boundary invest in custom spelling and aliases, while short-lived pipeline dialects fall back on the generic printer. Expect polished textual forms for cuda_tile, cute, and cute_nvgpu; expect generic MLIR for most nv_tileaa, nv_tileas, and cutlass operations, with a few aliases or SSA-name hints sprinkled in to keep large dumps legible.

Alias resolution

OpAsmDialectInterface exposes two virtual hooks the printer consults before falling back on the generic form: getAlias(Type, raw_ostream&) and getAlias(Attribute, raw_ostream&). Each hook returns an AliasResult — NoAlias, OverridableAlias, or FinalAlias — and, when the result is not NoAlias, writes the alias name into the stream. The printer queries every loaded dialect in registration order; the first non-NoAlias answer wins. FinalAlias short-circuits subsequent dialects; OverridableAlias permits a later dialect to refine the name.

bool emit_type_with_alias(AsmPrinter *p, Type t) {
    SmallString<32> name;
    raw_svector_ostream os(name);
    for (Dialect *d : p->context()->loadedDialects()) {
        OpAsmDialectInterface *iface = d->getRegisteredInterface<OpAsmDialectInterface>();
        if (!iface) continue;
        AliasResult r = iface->getAlias(t, os);
        if (r == AliasResult::NoAlias) { name.clear(); continue; }
        register_alias_decl(p, t, name);   /* emit `!name = type ...` at top of module */
        p->os() << "!" << name;
        return true;
    }
    return false;  /* caller falls back to generic !ns<...> form */
}

When emit_type_with_alias returns false the printer writes the generic form — !ns<storage-blob> for parametric types, the registered mnemonic plus storage for ODS-generated types. Attribute printing follows the same shape with # in place of !.

Per-dialect feature matrix

The table below summarizes which textual-IR hook each dialect installs. "ODS-only" means the slot is wired by the TableGen-generated dialect registration to MLIR's default trampoline (which reads the registered mnemonic/storage and emits the canonical form). "stub" means the slot is patched to a body that either does nothing or emits a parsing in dialect '<ns>' is disabled diagnostic. "real" means a hand-written dispatcher of non-trivial size.

cuda_tile — user-facing input syntax

cuda_tile has the richest textual surface. Constants receive stable SSA-name hints — cst, true, false, cst_NaN, cst_<int> — that keep debug dumps legible. Selected TKO load/store and atomic operations carry hand-written printers and parsers instead of generic MLIR spelling.

nv_tileaa — generic dialect with a few name hints

nv_tileaa installs no dialect-wide asm aliases — most operations print in generic MLIR form. Six operations attach per-op asm interfaces, and the only pretty-name behavior worth knowing lives on nv_tileaa.load: the value result is named result, and the optional memory-token result is named resultMemToken.

nv_tileas — aliases for scheduling concepts

nv_tileas falls back on generic operation printing for most ops but ships useful dialect-level aliases for scheduling attributes and types. Attribute aliases cover memory-space layouts, copy atoms, reduction atoms, MMA atoms, and resource requirements. Type aliases cover pipeline and role-qualified iterator types such as producer and consumer iterators.

cute — attributes are the serialized type surface

cute disables standalone type parsing. Its canonical textual form represents types as #cute.<keyword> attributes rather than !cute.<keyword> types. The attribute parser recognizes layout-algebra terms — coord, stride, shape, tile, swizzle, layout, composed_layout, ptr, memref, coord_tensor — along with constrained integer forms.

cute_nvgpu — architecture atom spelling

cute_nvgpu ships a full type parser/printer for architecture-specific MMA, copy, TMA, shared-memory descriptor, and tensor-memory atoms. Its OpAsmDialectInterface aliases collapse the otherwise unwieldy stored type body — element-type triples, operand orderings, instruction shapes, swizzle patterns — into a short label that survives in 10 000-line dumps.

The alias hook is a discriminated dispatch on the type's class id. Each branch builds the alias name from the type's own accessors rather than from the encoded storage body, so the alias survives parameter reordering inside the storage struct.

AliasResult cute_nvgpu_type_alias(Type t, raw_ostream &os) {
    if (auto m = dyn_cast<MemRefAtomType>(t)) {
        os << "memref_" << element_type_keyword(m.getElementType())
           << "_"      << m.getRank();
        return AliasResult::OverridableAlias;
    }
    if (auto c = dyn_cast<CopyAtomType>(t)) {
        os << "Cp(" << element_type_keyword(c.getElementType()) << ","
                    << c.getShape().M << "x" << c.getShape().N << ","
                    << layout_keyword(c.getLayout()) << ")";
        return AliasResult::FinalAlias;
    }
    if (auto m = dyn_cast<MmaAtomType>(t)) {
        os << "Mma(m" << m.getInstShape().M
           <<   "n"  << m.getInstShape().N
           <<   "k"  << m.getInstShape().K << ","
           << element_type_keyword(m.getElementTypeA()) << ","
           << layout_keyword(m.getLayoutA()) << ","
           << layout_keyword(m.getLayoutB()) << ")";
        return AliasResult::FinalAlias;
    }
    if (isa<TmaDescriptorType, SharedDescriptorType, TensorMemoryType>(t))
        return tma_family_alias(t, os);
    return AliasResult::NoAlias;
}

The MMA alias exposes the instruction shape and element-type/layout triple without expanding the full atom storage body. A printed cute_nvgpu dump that would otherwise contain !cute_nvgpu.mma_atom<inst_shape = <m = 16, n = 8, k = 16>, a = <element = f16, layout = row>, b = <element = f16, layout = col>, c = <element = f32>> instead reads !Mma(m16n8k16,f16,row,col). The 23 typed copy and MMA atoms collapse to one short label each.

cutlass — generic spelling

cutlass leaves textual assembly to the framework on purpose. Its operations carry registered attributes and opaque types, so the generic ODS printer produces sufficient IR without dialect-wide aliases or per-op pretty names.

Practical rules

A reimplementer who wants compatible textual dumps starts at the cuda_tile boundary because every input module passes through it: stable constant SSA hints (cst, true, false, cst_NaN, cst_<int>) and the hand-written TKO load/store/atomic printers carry the largest readability payoff. The nv_tileaa.load result names result and resultMemToken are fixed contract — downstream regression dumps reference them. The cute invariant that all type-like syntax appears as #cute.<keyword> attributes rather than !cute.<keyword> types must be preserved because the textual parser refuses standalone cute types outright. The cute_nvgpu aliases for memref, copy, and MMA atoms cover the bulk of the alias-driven readability gain; per-atom custom printers can come later. Operations in cutlass and most of nv_tileas ship without aliases on purpose — the generic ODS form is precise enough and avoids divergence between the textual dialect description and the dialect's stored representation.

Cross-references

The companion page Dialect Bytecode Reader/Writer Status covers the input-side wire format that produces the IR these printers later spell. Layout-algebra spelling for the cute attribute family is documented in cute Layout Algebra and Descriptor Grammar; tensor-memory and TMA atom encodings are described under cute_nvgpu Asm Printer and Mnemonic Hash and cute_nvgpu TMA Atoms. The architectural rationale for keeping intermediate dialects on the generic textual form lives in cuda_tile Asm Printer.

cuda_tile Dialect Overview

Frontends write cuda_tile and the compiler promises to accept it. It is the public input contract of tileiras — the only dialect a producer ever has to construct — and the gate before lowering descends into the private TileAA, TileAS, CuTe, CUTLASS, NVGPU, LLVM, and NVVM layers. In practice it is a compact tile-programming IR: structured control flow, shaped tile values, view-based memory access, token-threaded side effects, tensor-core operations, and just enough attributes to preserve numeric and memory semantics until target-specific lowering takes over.

Producers generate cuda_tile; reimplementers treat it as an ABI boundary. A module that verifies here flows through the rest of the compiler without the frontend ever touching nv_tileaa, nv_tileas, or any backend dialect.

Programming Model

A normal input module is rooted in cuda_tile.module and contains one or more cuda_tile.entry operations that each become a GPU kernel. Inside each entry, the dialect carries its own structured control flow (if, for, loop, yield, break, continue, return) so frontends never have to lower into scf or func first.

Values fall into four broad categories:

The dialect is target-aware but not target-lowered. Accepted element types are f16, bf16, f32, tf32, f64, f8E4M3FN, f8E5M2, and the integer widths i1, i8, i16, i32, i64. Architecture-specific choices — MMA atom selection, TMA materialization, register allocation, FP4/FP6 microscaling, final PTX features — all come later, in the private lowering pipeline.

Operation Families

The operation surface is best understood by family rather than by registration order:

The exact roster is maintained in Operation Roster. Two practical version deltas matter for producers targeting this binary: the emitted mnemonic is cuda_tile.print, not the open-source cuda_tile.print_tko, and the build rejects cuda_tile.atan2 outright.

Type Contracts

cuda_tile types describe the source-level shape and memory model. They should be treated as verifier-backed contracts, not as backend storage layouts.

The tile-shape verifier walks the shape, rejecting non-positive or non-power-of-two dimensions and enforcing a 16-million-element ceiling. The element count is tracked using a divide-and-compare to detect overflow before it can happen.

LogicalResult verify_tile_shape(ArrayRef<int64_t> shape) {
    const int64_t max_elements = 16 * 1024 * 1024;
    int64_t elements = 1;

    for (int64_t dim : shape) {
        if (dim <= 0) {
            return emit_error("tile dimensions must be positive");
        }
        if ((dim & (dim - 1)) != 0) {
            return emit_error("tile dimensions must be powers of two");
        }
        if (elements > max_elements / dim) {
            return emit_error("tile would exceed the maximum element count");
        }
        elements *= dim;
    }
    return success();
}

tensor_view uses dynamic shape and stride slots, but each dynamic slot is still part of a fixed-rank type. The verifier rejects rank mismatches between shape and stride and rejects any non-positive static dimension or static stride.

LogicalResult verify_tensor_view(Type element_type,
                                 ArrayRef<int64_t> shape,
                                 ArrayRef<int64_t> stride) {
    if (shape.size() != stride.size()) {
        return emit_error("tensor_view shape and stride must have the same rank");
    }
    for (int64_t dim : shape) {
        if (dim != kDynamic && dim <= 0) {
            return emit_error("static tensor_view dimensions must be positive");
        }
    }
    for (int64_t step : stride) {
        if (step != kDynamic && step <= 0) {
            return emit_error("static tensor_view strides must be positive");
        }
    }
    return success();
}

partition_view is the bridge between logical tensors and tile-shaped access. The verifier checks rank agreement, validates dim_map as an injective function into tensor axes, enforces the power-of-two tile shape rule, and gates special padding values on floating-point element types.

LogicalResult verify_partition_view(ArrayRef<int32_t> tile_shape,
                                    TensorViewType tensor,
                                    ArrayRef<int32_t> dim_map,
                                    Optional<PaddingValue> padding) {
    if (tile_shape.empty()) {
        return emit_error("partition tiles must have rank");
    }
    if (tile_shape.size() != tensor.rank()) {
        return emit_error("partition tile rank must match tensor rank");
    }
    if (dim_map.size() != tile_shape.size()) {
        return emit_error("dim_map must cover every tile dimension");
    }

    BitSet used_tensor_dims(tensor.rank());
    for (size_t tile_dim = 0; tile_dim < dim_map.size(); ++tile_dim) {
        if (tile_shape[tile_dim] <= 0) {
            return emit_error("partition tile dimensions must be positive");
        }
        if (!is_power_of_two(tile_shape[tile_dim])) {
            return emit_error("partition tile dimensions must be powers of two");
        }
        int32_t tensor_dim = dim_map[tile_dim];
        if (tensor_dim < 0 || tensor_dim >= (int32_t)tensor.rank()) {
            return emit_error("dim_map target must be inside the tensor rank");
        }
        if (used_tensor_dims.test(tensor_dim)) {
            return emit_error("dim_map must not map two tile dimensions to one tensor dimension");
        }
        used_tensor_dims.set(tensor_dim);
    }

    if (padding.has_value() && padding->is_nan_or_infinity_or_negative_zero()) {
        if (!tensor.element_type().is_float()) {
            return emit_error("special padding values require a floating-point element type");
        }
    }
    return success();
}

Memory and Tokens

The _tko suffix means token-ordered. Memory effects ride on dataflow: the token is an SSA value, and a pass may reorder memory operations only when it preserves the dependency graph that ties them together.

struct Token {};

struct LoadResult {
    Value value;
    Token token;
};

LoadResult load_ptr_tko(Pointer ptr, Indices indices, Token in);
LoadResult load_view_tko(PartitionView view, Indices indices, Token in);

Token store_ptr_tko(Pointer ptr, Indices indices, Value value, Token in);
Token store_view_tko(PartitionView view, Indices indices, Value value, Token in);

struct AtomicResult {
    Value old_or_result;
    Token token;
};

AtomicResult atomic_rmw_tko(Pointer ptr, AtomicOp op, Value value, Token in);
AtomicResult atomic_cas_tko(Pointer ptr, Value expected, Value desired, Token in);

A pass may delete, merge, or reorder token-ordered operations only when the observable token order survives intact. That is the source-level memory contract that later TileAA and TileAS passes refine into schedulable memory operations.

Semantic Attributes

The attribute set is small but consequential. Most attributes are not decoration — they constrain legal lowering:

Key design choice: public because it's the API

cuda_tile is public because it is the producer-facing API. Every dialect below it is an implementation detail. A frontend should construct valid cuda_tile, serialize it as TileIR bytecode, and hand it to tileiras — never touching internal TileAA or TileAS operations.

The lowering direction is one-way. The driver runs in three phases.

Phase 1: the verifier rejects modules that contain operations from any non-public dialect. A producer that emits IR through this entry point must restrict itself to cuda_tile, builtin, and a small set of supporting upstream dialects (arith constants, func symbol references, debug-info attributes). Any other dialect at this point is a producer bug.

Phase 2: a partial dialect conversion drives the rewrite. The conversion target marks cuda_tile illegal, marks the destination dialects (arith, math, func, gpu, scf, nv_tileaa) legal, and registers a dynamic legality check on ub.poison so untyped poison values pick up legal TileAA types as they flow through. Each cuda_tile op carries a conversion pattern that emits the corresponding TileAA shape; the type converter rewrites scalar, tile, pointer, view, and token types in parallel.

Phase 3: a post-conversion verifier confirms that no cuda_tile operation survived the conversion. After this point, ordinary producers will never see cuda_tile again; the rest of the pipeline works in progressively more hardware-facing internal dialects (TileAA → TileAS → CuTe → NVGPU → NVVM → LLVM).

The driver is structured as a single greedy pass rather than a per-family sweep because the rewrite patterns produce IR that immediately matches further patterns: a cuda_tile.load_view_tko lowers into a TileAA tiled_load that exposes new shape and layout structure to the next op's lowering. A per-family sweep would force a fixed phase order; the greedy pass lets pattern match order respond to the IR as the conversion produces it.

Open-source cross-reference

The public cuda_tile source distribution is the best reference for syntax, ODS definitions, operation classes, type definitions, and dialect interfaces. The binary follows that public surface with the practical deltas noted above: print_tko is exposed as print, atan2 is absent, and this binary also contains an implementation-specific cuda_tile.string type.

The useful public source anchors are:

AbstractOperation Record

Every registered op in cuda_tile carries one AbstractOperation descriptor. The dialect constructor walks its 92-op roster, allocates one descriptor per op, fills it from that op's registration thunk, and appends it to the dialect's registered-op vector. An Operation* resolves through its OperationName slot into this descriptor to reach the dialect's interface tables and fold callback.

The descriptor's logical layout:

The descriptor slab is zero-initialized, so unused interface slots stay null and the dispatcher probes them without a presence flag. The mnemonic field is an embedded StringRef that points at the binary's read-only literal, not a heap-interned copy — the ASM printer and the verifier read it back verbatim.

The descriptors sit consecutively in a statically-allocated array. The dialect indexes the array by mnemonic hash through the registration helper documented in TypeID Sentinels and Anchors; live Operation* instances reach the descriptor through their OperationName slot — the resolution path documented in Operation Layout — Pointer-Identity Dispatch. The per-op fold-callback assignments for the rest of the roster are catalogued in Operation Roster — Op Method Surface.

Cross-links

• Frontend Contract and Tile IR Emission — producer-facing rules for kernel signatures, attribute namespaces, operand- order conventions, and the bytecode-format constraints a conformant frontend must satisfy.
• Operation Roster — operation families, producer contract, and version-specific mnemonic notes.
• Types and Attributes — public types, element predicates, semantic attributes, assumption predicates, and optimization hints.
• Verifiers — numeric, memory, region, aggregate, and MMA verification contracts.
• Canonicalizers and Folds — public folds, select and if rewrites, and the recursive simplifier contract.
• Assembly Printer — textual assembly, token-memory syntax, attribute elision, enum spellings, and SSA result-name hints.

cuda_tile Operation Roster

Abstract

A frontend emitting cuda_tile is writing tile values, structured kernel control flow, token-ordered memory effects, tensor views, matrix multiply-accumulate intent, and source-level numeric attributes — everything the compiler will subsequently lower into private implementation dialects. This page is the producer and reimplementation reference: operation families, the behavior each family promises, and how a compiler should lower the surface without leaning on internal registration details.

In this build, the token-ordered print operation is spelled cuda_tile.print. The newer cuda_tile.atan2 is rejected outright, so a frontend that supports multiple TileIR revisions should gate it behind explicit version logic.

Operation Families

The family boundaries are semantic, not syntactic. fma is arithmetic because it is elementwise; mmaf and mmai are MMA because they contract matrix dimensions. assert and assume live with control flow because regions and dominance scope their meaning, even though their payload is an attribute or predicate.

Producer Contract

A valid producer should build modules with this shape:

cuda_tile.module {
    cuda_tile.entry @kernel(%arg0 : !cuda_tile.tensor_view<...>) {
        %tok0 = cuda_tile.make_token : !cuda_tile.token
        %tile, %tok1 = cuda_tile.load_view_tko %view[%i, %j] token=%tok0
        %acc = cuda_tile.mmaf %a, %b, %c : ...
        %tok2 = cuda_tile.store_view_tko %view[%i, %j], %acc token=%tok1
        cuda_tile.return
    }
}

The exact textual syntax is described in Assembly Printer, but the contract is independent of formatting:

• memory effects are threaded through cuda_tile.token;
• tile values have static rank and element type;
• view values carry shape and stride metadata;
• structured control flow yields values rather than branching through cf;
• numeric choices such as rounding and signedness are attributes, not implicit frontend assumptions;
• debug info and optimization hints may be present but must not be required for semantic correctness.

Lowering Sketch

The first lowering stage converts public cuda_tile into alias-aware TileAA. Arithmetic and shape operations keep their mathematical meaning intact. Memory operations gain explicit memref and token structure. Control flow is rewritten only once region and token legality are already proven.

Module lower_cuda_tile_to_tileaa(Module module, Target target) {
    require(module.only_uses_dialect("cuda_tile", "builtin", "arith"));
    verify_cuda_tile_module(module, target);

    TypeConverter types;
    types.add(convert_scalar_type);
    types.add(convert_tile_type);
    types.add(convert_pointer_type);
    types.add(convert_view_type);
    types.add(convert_token_type);

    RewritePatternSet patterns;
    add_arithmetic_patterns(patterns, types);
    add_shape_patterns(patterns, types);
    add_memory_patterns(patterns, types);
    add_control_flow_patterns(patterns, types);
    add_mma_patterns(patterns, types);

    apply_conversion(module, patterns);
    require(!module.contains_dialect("cuda_tile"));
    return module;
}

Lowering must not erase source-level facts prematurely. A load_view_tko becomes an operation with explicit view, index, mask, fallback, memory ordering, memory scope, and token dependencies — not an unstructured pointer load until the alias and layout passes have the context to handle it safely.

Numeric Operations

Arithmetic ops accept scalar or tile-shaped operands. Tile operands must agree on shape and element type unless the op has an explicit shape-changing contract. Floating operations carry rounding mode and flush-to-zero policy forward until a lower dialect decides whether the target instruction can encode those choices directly.

Value lower_elementwise_arith(ArithOp op) {
    require_same_shape(op.operands);
    require_legal_element_type(op);

    NumericPolicy policy = {
        .rounding = op.rounding_mode,
        .flush_to_zero = op.flush_to_zero,
        .signedness = op.signedness,
        .overflow = op.overflow,
    };

    return tileaa_elementwise(op.kind, op.operands, policy);
}

mulhii returns the high half of a signed integer product. Implement it as a wide multiply followed by a high-half extract — never as ordinary multiplication that relies on target-width overflow.

Operand and Result Tables

The most heavily emitted ops carry the following operand/attribute/result shape. The _tko family threads a cuda_tile.token through every memory effect.

cuda_tile.load_view_tko

A representative two-dimensional tile load with a predicate mask:

%tile, %t1 = cuda_tile.load_view_tko %view[%i, %j], %mask, %fallback, %t0
    { mem_semantic = #cuda_tile<mem_semantic relaxed>,
      mem_scope    = #cuda_tile<mem_scope gpu>,
      operandSegmentSizes = array<i32: 1, 2, 1, 1, 1> }
    : !cuda_tile.partition_view<128x64xf32>, index, index,
      tile<128x64xi1>, tile<128x64xf32>, !cuda_tile.token
    -> tile<128x64xf32>, !cuda_tile.token

The mask and fallback shapes equal the result tile shape; the view element type matches the result element type. The token chain threads %t0 in and %t1 out, ordering the load against any preceding or following memory effect that consumes the same chain.

cuda_tile.store_view_tko

cuda_tile.atomic_rmw_tko

cuda_tile.mmaf / cuda_tile.mmai

A 16×16×16 floating MMA with an f32 accumulator and f16 inputs:

%d = cuda_tile.mmaf %a, %b, %c
    : tile<16x16xf16>, tile<16x16xf16>, tile<16x16xf32>
    -> tile<16x16xf32>

A batched integer MMA with explicit signedness attributes:

%d = cuda_tile.mmai %a, %b, %c
    { signedness_a = #cuda_tile<signedness signed>,
      signedness_b = #cuda_tile<signedness unsigned> }
    : tile<4x16x32xi8>, tile<4x32x16xi8>, tile<4x16x16xi32>
    -> tile<4x16x16xi32>

The M/N dimensions of A and B agree with C; the K dimension is contracted. The verifier rejects rank mismatch, K disagreement, accumulator/result type mismatch, missing signedness on mmai, and any input/accumulator pair that lies outside the target's legal MMA element-type tuple.

cuda_tile.if

cuda_tile.for

Memory and Token Operations

The _tko suffix means token ordered. Every token-ordered memory op consumes an input token and produces a successor. Loads and atomics also produce data; stores produce only the successor token. That discipline is the public memory model — later passes refine it into barriers, async copies, and backend memory instructions.

LoadResult lower_load_ptr_tko(LoadPtrTkoOp op) {
    MemRef ref = make_memref_from_pointer(op.pointer, op.indices);
    MemoryPolicy policy = memory_policy(op.ordering, op.scope, op.hints);

    Value data = tileaa_load(ref, op.mask, op.padding, policy, op.input_token);
    Token next = token_after(data.memory_effect, op.input_token);
    return (LoadResult){ .value = data, .token = next };
}

Atomics check both memory ordering and element type. Integer bitwise modes are integer-only; floating add is floating-only; compare-and-swap is restricted to element widths the backend can update atomically.

Structured Control Flow

cuda_tile ships its own region operations because frontends need a stable kernel-level API. Later lowering may translate these regions into SCF, CFG, or private control-flow dialects, but the verifier enforces these rules first:

• if result types match every non-empty yielding branch;
• for induction, bounds, step, iter args, and results are type-consistent;
• loop iter args and results are type-consistent;
• break exits the nearest compatible loop;
• continue exits to the next iteration of a compatible for or loop;
• return appears in an entry context and matches the entry function type;
• yield appears only in a parent op that expects region yields.

MMA Operations

mmaf and mmai are deliberately narrow public abstractions: they describe matrix multiply-accumulate intent, not final tensor-core instruction selection. The verifier checks shape compatibility and element-type legality. Choosing WGMMA, smaller MMA atoms, tensor-memory paths, or emulation is left to the lowering pipeline.

LogicalResult verify_mma_shape(Tile lhs, Tile rhs, Tile acc, Tile result) {
    require(lhs.rank == 2 || lhs.rank == 3);
    require(rhs.rank == lhs.rank);
    require(acc.rank == lhs.rank);
    require(result.rank == lhs.rank);

    if (lhs.rank == 3) {
        require(lhs.dim(0) == rhs.dim(0));
        require(lhs.dim(0) == acc.dim(0));
        require(lhs.dim(0) == result.dim(0));
    }

    require(lhs.k_dim == rhs.k_dim);
    require(lhs.m_dim == acc.m_dim);
    require(rhs.n_dim == acc.n_dim);
    require(acc.shape == result.shape);
    return success();
}

Version Notes

• Emit cuda_tile.print for runtime diagnostic printing in this build.
• Do not emit cuda_tile.print_tko unless targeting a source tree that uses that mnemonic.
• Do not emit cuda_tile.atan2 for this build; guard it behind a newer TileIR version check.
• Treat cuda_tile.string as implementation-specific unless the target contract explicitly documents it.

Op Method Surface

Every cuda_tile.* op exposes four registered functions to the framework: a builder (or textual parse entry), a registration thunk that interns the mnemonic and installs the op vtable, a verifier hook, and a lowering pattern that rewrites the op during the first conversion stage. The functions follow predictable shapes by op family.

The default builder for trivial unary ops shares one trampoline that constructs the op from a single operand and forwards rounding/flush-to-zero attributes; the default verifier hook installs a no-op stub when the op's contract is fully covered by trait-level checks. The control-flow lowering routes through one driver that owns cuda_tile.if, cuda_tile.for, cuda_tile.loop, cuda_tile.continue, and cuda_tile.return together so it can preserve region nesting and the structured-exit ancestry contract.

One count discrepancy is worth flagging. The roster in this build is 92 mnemonics. The two names missing from open-source documentation are cuda_tile.atan2 (excluded entirely from this binary) and the rename cuda_tile.print_tko → cuda_tile.print. Producers should follow the version notes above and emit only the 92 mnemonics this dialect accepts.

The dialect constructor walks the registration thunks in roster order; each thunk interns the mnemonic into the dialect's OperationName table and installs the op's vtable, fold callback, and verifier hook through the slots described in overview — AbstractOperation Record and Operation Layout — Pointer-Identity Dispatch. Lowering patterns are matched as conversion patterns by the arithmetic-group and pointer-cast dispatchers during the first lowering stage; the conversion is documented in Cuda Tile to TileAA.

Cross-References

Overview describes the dialect's role as the public producer-facing API and the AbstractOperation record structure. Verifiers details the verbatim verifier diagnostics each family emits. Canonicalizers and Folds describes the rewrites applied after verification. Bytecode Reader and Writer documents the on-wire encoding the opcode dispatcher consumes.

cuda_tile Types and Attributes

Abstract

cuda_tile types are the public shape and memory vocabulary of TileIR: tile values, typed pointers, tensor views, partitioned views, and ordering tokens. Attributes layer on the numeric, memory, target, padding, assumption, optimization, and debug facts that make lowering deterministic. This page lays out those contracts in the terms a frontend or reimplementation needs — which values may be constructed, which attributes are semantic, and which facts are verified before the module enters private lowering.

Concrete Types

The first five types form the stable public surface. cuda_tile.string is useful for reading this build's dumps; portable producers must not depend on it.

TypeStorage and TypeID Singletons

Each of the six concrete types is a normal MLIR Type subclass backed by its own TypeStorage derivative. Construction flows through the MLIRContext's StorageUniquer gateway documented in Storage Uniquer and Context Impl — getOrCreate Gateway: the dialect's self-registration ctor hands a TypeID singleton and a build hook to the uniquer, which hashes the storage payload, looks up an existing instance, and either returns the cached pointer or allocates a fresh storage block in the context arena. Every public Type handle in the IR is a 24-bit-tagged pointer into that arena.

All six types share a 24-byte BaseStorage header (vtable, context pointer, hash-bucket pointer) and then append a type-specific payload. Storage sizes are byte-exact across the build and stable under bytecode round-trip.

The PointerType singleton &unk_5B38BC8 is the exact value the TileElementType predicate (sub_6C4E20) tests against in its final arm when it accepts a typed pointer as a tile element — the registration ctor's TypeID slot is observably the same one driving tile-element verification.

TileType storage

cuda_tile.tile carries a static shape and an element type. The shape is held as an ArrayRef<int64_t> (begin pointer plus size, 16 bytes); the element type is a single tagged pointer.

typedef struct TileTypeStorage {
    /*+0x00*/ BaseStorage    base;          // vtable=&unk_5B38BC0, ctx, hash bucket
    /*+0x18*/ const int64_t *shape_begin;   // pointer into context-owned int64 array
    /*+0x20*/ uint64_t       shape_size;    // dimension count
    /*+0x28*/ Type           element_type;  // scalar element or cuda_tile.ptr
} TileTypeStorage;

The shape array is interned alongside the storage block, and copies returned to callers re-use that pointer. The element type field accepts the TileElementType palette — f16, bf16, f32, tf32, f64, f8E4M3FN, f8E5M2, i1, i8, i16, i32, i64, and cuda_tile.ptr. Total storage 0x30 bytes.

TensorViewType storage

cuda_tile.tensor_view carries an element type, a shape, and a stride. Both shape and stride are full ArrayRef<int64_t> records, and both accept the shared MLIR kDynamic = INT64_MIN sentinel independently per dimension.

typedef struct TensorViewTypeStorage {
    /*+0x00*/ BaseStorage    base;          // vtable=&unk_5B38BB8
    /*+0x18*/ Type           element_type;
    /*+0x20*/ const int64_t *shape_begin;
    /*+0x28*/ uint64_t       shape_size;
    /*+0x30*/ const int64_t *stride_begin;
    /*+0x38*/ uint64_t       stride_size;
} TensorViewTypeStorage;

Element type must satisfy the bare-Number predicate (sub_6C2A10): the tile-element palette minus the cuda_tile.ptr arm. Total storage 0x40 bytes.

PartitionViewType storage

cuda_tile.partition_view overlays a power-of-two tile grid on a tensor view and optionally selects a padding value for out-of-bounds reads. The tile shape is held as a DenseI32ArrayAttr attribute pointer (interned independently by the attribute uniquer); the dimension map is a raw ArrayRef<int32_t>; the padding value is a nullable attribute slot.

typedef struct PartitionViewTypeStorage {
    /*+0x00*/ BaseStorage         base;           // vtable=&unk_5B38BB0
    /*+0x18*/ DenseI32ArrayAttr   tile_shape;     // interned attr
    /*+0x20*/ TensorViewType      tensor_view;
    /*+0x28*/ const int32_t      *dim_map_begin;
    /*+0x30*/ uint64_t            dim_map_size;
    /*+0x38*/ PaddingValueAttr    padding_value;  // nullable, attr pointer
} PartitionViewTypeStorage;

At registration time, the self-registration ctor also wires the TileView interface concept-model pointer and method table into the TypeStorage+0x88 slot. Op verifiers such as verifyViewLoadStoreCommon reach the partition view through that interface to read getViewIndexRank() and getViewTileType(), so the interface vtable is consulted on every view load and store. Total storage 0x40 bytes.

PointerType storage

cuda_tile.ptr is a typed pointer to a single scalar element. This build carries no explicit address-space field: the pointer is always a global-memory typed reference, and any address-space variation rides on the tensor_view or partitioned access it flows through, not on the pointer type itself.

typedef struct PointerTypeStorage {
    /*+0x00*/ BaseStorage    base;          // vtable=&unk_5B38BC8
    /*+0x18*/ Type           pointee_type;  // bare-Number palette only
} PointerTypeStorage;

The pointee type must satisfy the bare-Number predicate (sub_6C2840); pointer-to-pointer is forbidden by that arm. Total storage 0x20 bytes.

TokenType storage

cuda_tile.token is parameter-free. It carries no payload beyond the shared BaseStorage header — its only job is to thread ordering edges between token-producing and token-consuming memory operations.

typedef struct TokenTypeStorage {
    /*+0x00*/ BaseStorage    base;          // vtable from off_5A2E208 slot
} TokenTypeStorage;

Total storage 0x18 bytes. Two cuda_tile.token SSA values produced by different ops are unequal IR values, but their storage instance is unique — every !cuda_tile.token type in a context resolves to the same TypeStorage.

StringType storage

cuda_tile.string is also parameter-free at the storage layer in this build. It is an internal handle used by debug and diagnostic plumbing; public producers should treat it as nonportable.

typedef struct StringTypeStorage {
    /*+0x00*/ BaseStorage    base;          // vtable from off_5A2DB38 slot
} StringTypeStorage;

Total storage 0x18 bytes. Like cuda_tile.token, the type resolves to a single canonical storage instance per context.

Element-type dispatch and the 11-arm predicate

The TileType self-registration ctor at sub_6C5870 wires the TileElementType predicate at sub_6C4E20 into the parameter-trait verifier emitted by TableGen. That predicate is an unrolled AnyTypeOf<> switch over the element-type singleton table — each arm a direct vtable-pointer test or a width-keyed isInteger(N) call. The accepted set has thirteen arms in the binary (f16, bf16, f32, tf32, f64, f8E4M3FN, f8E5M2, i1, i8, i16, i32, i64, and cuda_tile.ptr) and emits the verbatim failure string failed to verify 'elementType': f16 or bf16 or f32 or tf32 or f64 or f8E4M3FN or f8E5M2 or i1 or i8 or i16 or i32 or i64 or Pointer type on mismatch. The bare-Number predicate at sub_6C2A10 is the same dispatch table minus the final cuda_tile.ptr arm — which is how the verifier forces tensor_view element types to be scalar while still letting pointer elements live inside tiles.