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.