TileFoundry Spec — Code Organization

Implementation guide (not an architecture spec). Defines the Python source tree layout.

1. Directory skeleton

This spec describes the stable layers only — whether a single .py file exists is decided by the naming rules in §2 and is not enumerated here (so adding a new Op does not require a spec edit). The current file list is whatever git ls-files src/tilefoundry reports.

The top-level package src/tilefoundry/ is divided as follows. Each directory has an owning spec — that spec is the single source of truth for the directory's structure and invariants.

Directory	Owning spec	Contents
ir/core/	core-ir	Shared node algebra: Module / Expr / Var / Constant / Tuple / Op / Call / Stmt (base class) / OpSchema / ParamDef / definition call-path enumeration / typed metadata attach-detach and diagnostics / @register_op / @register_alias / op_registry / errors.
ir/pattern/	core-ir	Operation-declaration predicates: composable pattern values in pattern.py, match/binding and rendering mechanics in match.py, cross-operand relations in constraint.py, and pattern construction/specialization helpers in utils.py.
ir/types/	types	Type-system root: Type / TensorType / TupleType / UnitType / CallableType / DType / StorageKind / resolve_storage / local projections (local_type_of) / tensor-leaf, byte-by-storage, and topology-extent queries / dim.* (with their typeinfer).
ir/types/{int_tuple,stride,layout,layout_algebra,swizzle_layout,shard_layout,mesh}.py	shard	Topology / Mesh / Layout / ComposedLayout / ShardLayout / ShardAttr (Split / Broadcast / Dynamic / Partial), filed as CuTe files them: int tuples (flatten / unflatten / repeat_like / product), strides (compact_major / idx2crd / crd2idx), layouts, general layout algebra, and swizzle-specialized layout algebra each in their own module; mesh construction, separation, and topology-bound checking stay with Mesh.
ir/mesh_scope.py	shard	Which scope a statement stands inside and what it admits: device_layout, covered_by_scope. Neither a type nor a visitor, so it sits beside ir/isl_interop.py rather than in either.
ir/clause/	parser	Authored where(layout=..., mesh=..., storage=...) constraint records: the shared base plus layout, mesh, and storage constraints, attached by the parser and read back by the Python printer.
ir/visitor.py	visitor-mutator	ExprFunctor / ExprVisitor / ExprWalker / ExprCollector / ExprCloner / BindingSubstitutionCloner / StmtVisitor / StmtMutator / StmtExprMutator, plus collect_exprs, value-operand/function-value queries, and the canonical PrimFunction walk and rewrite entries.
ir/isl_interop.py	types §11	Interoperation between dimension and shape IR values and isl: expression rendering and decoding, normalization, value ranges, and shape-domain construction. Pure isl operations remain in utils/isl_utils.py.
ir/hir/	hir	HIR Op layer; one subdirectory per multi-operation category (math/ / tensor/ / nn/ / shape/ / sharding/), with the standalone scheduling carrier in schedule.py. One real Op per .py (§2 rule 1); surface-alias schemas have no per-name file and live in each category's aliases.py (§2 rule 5).
schedule/	schedule	Schedule finalization, instruction facts, and candidate discovery over HIR.
ir/tir/	tir	TIR layer: stmt.py re-exports the Stmt base from ir/core/stmt.py; stmts.py hosts the general TIR Stmt subclasses (LetStmt / Evaluate / Sequential / MeshScope / …); prim_function.py; effect Ops and TIR-owned Expr Ops by category (memory/ / nn/ / …); launch.py owns Launch and its authored launch-attribute descriptors; arith.py / reduce.py for tag-dispatched Binary / Unary / Reduce; intrinsic.py for the @intrinsic decorator. Target-specific nodes nest under ir/tir/<target>/<category>/ (e.g. ir/tir/cuda/nn/mma.py) per §2 Rule 1c.
parser/	parser	DSL → IR parsing: base.py (shared visitor base + dispatch), hir_parser.py (@func body), tir_parser.py (@prim_func body), layout sugar / range-slice / dispatch modules. Not under ir/ — the parser is a producer of IR, not an IR sublayer.
analysis/	analysis	Fact layer over typed HIR: one module per analysis family. The compact public surface lives in analysis/__init__.py; per-target Facts projections live with their owning Target.
passes/	passes	Pass framework (pass_base.py / pass_manager.py) plus concrete transforms (transforms/<pass_name>.py, §2 rule 6).
target/	target	Compilation Target classes, class registration, service selection, and architecture/device facts: base.py owns Architecture / Device / Target / register_target / registered_targets; services.py owns the immutable service descriptors; each backend owns its concrete Target. Authored code constructs Target values; there is no string resolver.
target/cpu.py	target	The CpuTarget backend and its CPU code-generation service selection.
target/hardware/	target	The installed hardware database and its generic machinery: the authored Architecture / Device documents, the envelope and evidence-leaf loader, HardwareSpecRegistry, and the exact-key schema reader. It fixes the envelope only; the fact namespace below facts belongs to the target package named by a document's schema.
target/<backend>/spec.py	target	One backend's typed hardware schemas and the documents it installs, registered into the shared registry as an import side effect. This is where a fact path, its unit, and its cross-field invariants are validated, and where a document becomes an immutable Architecture / Device value.
analysis/api.py	analysis	The public composed Analyze operation: shared authored-program check and normalization, one per-call AnalyzeContext, dependency closure, ordering, single execution per member, Metadata-ownership enforcement, and semantic result assembly.
analysis/registry.py	analysis	The built-in Analyzer declarations. It re-exports the immutable Analyzer descriptor from target/services.py and holds no Target dispatch table.
analysis/errors.py	analysis	AnalysisError, the one diagnostic the whole analysis layer raises, so catching an analysis failure catches every analysis failure rather than the subset the caller happened to import.
analysis/visitor.py	analysis	The per-call AnalyzeContext, carrying the shared root/current IterationScope while a family traverses its work.
analysis/iteration_scope.py	analysis	The shared IterationScope tree built once from normalized HIR; families query it instead of constructing parallel structure.
analysis/access.py	analysis	Access relations resolved against authored iteration scopes, retaining each input's original boundary index.
analysis/loop_domain.py	analysis	isl iteration domains built from authored loop bounds.
analysis/loop_terms.py	analysis	IR-only resolution of constants, authored loop axes, strides, and bounded invariant offsets into LoopTerm; it has no isl dependency.
analysis/footprint.py	analysis	Reserved for future target-independent authored-loop footprint policy; access-relation construction stays in analysis/access.py.
analysis/report.py	analysis	Structured analysis report data, including record-family registration and field serialization. It depends only on analysis/core modules; inspection consumes it to produce text and source annotations.
analysis/check.py	analysis	The shared authored-program gate for analysis: authored-type re-derivation, authored validation, call-context validation, and checker-specific input checks. Established once per public call rather than per family. program_dim_vars publicly collects dimension declarations reached through program values and execution geometry.
analysis/facts.py	analysis	The narrow Facts aggregates the analysis families declare — the memory hierarchy graph, the throughput rates, and the parallel capacity. It is the record of how much hardware each measurement rests on, and names no backend; a Fact shared across consumer families belongs in target/facts.py.
analysis/metadata.py	analysis	The typed records the families leave on the IR, split by what each number depends on rather than by convenience.
analysis/compute_cost.py	analysis	The compute-cost family: logical flops per DType and bytes per storage level, from the authored program alone.
analysis/memory.py	analysis	The memory family: per-Call movement plus Function-wide movement, value lifetimes, per-level peaks, and non-fatal capacity errors for over-full addressable levels.
analysis/roofline.py	analysis	The roofline family: the recorded work divided by the target's published rates, per Call and aggregated per Function. Adds no count of its own.
analysis/performance.py	analysis	The performance family: occurrences projected from the shared IterationScope tree into flat timeline records and one function envelope, scaled by parallel capacity. It introduces no second scope tree.
visitor_registry/	visitor-registry	Shared registry instances and derived visitors: access-relation construction, contexts, ISL helpers, relation building, shard propagation, type inference, verification, code generation, and cost evaluation.
visitor_registry/op_cost.py	analysis	Each operation's per-instance flops and bytes, registered into the shared cost-evaluator registry. Owned here rather than by any target package, because the work an operation asks for follows from its own semantics and operand types on every backend.
inspection/analysis_report.py	inspection	Presentation of analysis-owned report data as text and annotated source. Analysis owns the structured report data and JSON dump; inspection owns how a human reads it.
target/<backend>/facts.py	target	One backend's Facts projections selected by its Target's get_facts. They restate installed documents in the shape a family declared and measure nothing.
target/facts.py	target	Facts used across consumer families, such as topology limits, plus validation for values returned by Target.get_facts: the requested frozen-dataclass shape and returned type. A Fact used by one family stays with that family; this module holds no projection registry.
codegen/	codegen	Code generation: the immutable CodeGenerator service; Linkable products in linkable_function.py / linkable_module.py; Linked products in linked_function.py / linked_module.py; linking in link_function.py / link_module.py; and concrete generators under <target>/ (mirroring ir/tir/ file layout — tir/<category>/<name>.py emitter, §2 rule 2). Not under ir/ — codegen is a consumer of IR; templates/ holds boilerplate only.
codegen/registry.py	codegen	The compatibility re-export of the immutable CodeGenerator descriptor and source-order grouping by equal Target values. It rejects multiple unequal CUDA Target groups before emission and holds no emitter registry.
runtime/	runtime	Runtime support (per-target headers, function templates, launch helpers).
inspection/	inspection	IR visualisation: DOT, Python printer, web viewer.
dump/	inspection	Dump flags, dynamically scoped dump contexts, and file/null dump sinks used by inspection and test integration.
dsl/	parser (authoring namespace)	User-facing import surface: tf/ (HIR namespace) / T/ (TIR namespace, including _platforms.py) / _namespace.py / _stub_gen.py / storage.py / __main__.py. The tf/__init__.pyi and T/__init__.pyi stubs are produced by python -m tilefoundry.dsl regen and are gitignored.
compile.py	architecture	tilefoundry.build / compile / jit top-level public verbs.
module.py	parser	The @module decorator entry point and module-level topology authoring constants.
script.py	parser	@func / @prim_func / @intrinsic decorator entry points.
__init__.py	inspection	The top-level view convenience entry and re-exports of public compiler surfaces.
utils/	code-organization	Shared leaf machinery: a module here MUST import nothing from ir/, parser/, passes/, codegen/, runtime/ or cli/, and MUST name no layer. It is depended on and depends on nothing, which is what lets a consumer outside the package — a pre-commit hook under an interpreter with nothing installed — load one of these modules by path and get the same implementation the package uses. A helper that needs to know a layer belongs in that layer; this is not a home for anything that did not fit.

Stage boundary. The pipeline picture in architecture §1 places parser/ and codegen/ outside ir/ (front-end producer and back-end consumer); the physical directory layout reflects that boundary directly.

Reading notes:

• ir/ holds the IR proper and its sublayers only. ir/types/ is the root of the type system, and the shard / layout types sit in it directly (architecture §3): they are types, and a folder naming them a sublayer said nothing a reader could use. Their filing follows CuTe's, one module per thing the algebra is written against.
• codegen/ and parser/ sit outside ir/. By the architecture §1 pipeline they are the front-end producer and back-end consumer of IR, not IR sublayers.
• analysis/ sits outside ir/ because it derives facts about typed HIR rather than defining an IR layer. It reads the IR and the Target and decides nothing over what it measures (architecture §5).
• codegen/<target>/ consumes TIR; an HIR Function enters only for skeleton emission. Its TIR subtree mirrors ir/tir/: prim_function lives in tir/, Stmt emitters in tir/stmts/, and memory/ / nn/ / arith/ / reduce/ / tensor/ each have their own subdirectory. codegen/<target>/hir/ holds only the skeleton emission of an HIR Function.
• Authored launch attributes belong to ir/tir/launch.py; the geometry a kernel is launched at is settled where the Launch is written (passes §7.3) and read back off it by codegen/topology.py, rather than derived a second time inside codegen or carried past it as a runtime-owned metadata type. The two launch contracts are distinct even though both are consumed across the codegen boundary.

ir/types/layout_algebra.py has a mechanically checkable public surface: every name in its __all__ MUST have the same name in CuTe, with is_inverse_projectable as its sole exception. Python's missing overload dispatch requires the separately named supports_composition, supports_inverse, and inverse adapters in swizzle_layout.py; NotProjectable in layout.py is the named Python diagnostic needed by composed application and inversion. These are explicit exceptions, not a license for unrelated helpers in the algebra modules.

ir/pattern/, ir/clause/, visitor_registry/, and dump/ are cross-cutting packages; their stable responsibilities are owned by core-ir, parser, visitor-registry, and inspection, respectively. Their internal file layout is not a per-Op contract. ir/pattern/ MUST NOT import inspection/, not even inside a function: a failed match answers with the matcher's Refusal, and the consumer that reports it turns that data into text with PatternPrinter.refusal. inspection/ may import ir/pattern/; the reverse edge does not exist. Likewise ir/tir/cuda/ atoms never import inspection/: an atom is a Printable from ir/core/, a printer hands itself to atom.print(printer, ctx), and the printer does not import an atom class.

2. File naming and content rules

Rule 1 — one real Op = one file. A real Op class lives in ir/<hir|tir>/<category>/<op_name>.py. The file name is the snake_case of the Op class CamelCase (MatMul → matmul.py, RMSNorm → rms_norm.py). TIR effect Ops and TIR-owned Expr Ops follow the same rule.

Rule 1a — surface-alias schemas have no per-name file. A surface alias (core-ir §2.3) has no IR class — its builder routes to a kinded target Op. All aliases for a category live together in aliases.py (e.g. the 19 HIR math sugar names add / sub / cmp_eq / neg / … all register in ir/hir/math/aliases.py).

Rule 1b — tag-dispatched IR classes. Binary / Unary / Reduce and other Op classes that fold many surface names through a kind attribute live in one file per IR class (ir/hir/math/binary.py / ir/hir/math/unary.py / ir/tir/arith.py / ir/tir/reduce.py). This does not contradict Rule 1: "one Op = one file" means one IR class per file; aliases are not IR classes, so they go through Rule 1a.

Rule 1c — target-specific IR nodes nest under the dialect. IR is dialect-first: its primary organizing axis is the dialect, and most nodes are target-neutral. A node or descriptor that is specific to one compilation target nests as ir/{dialect}/{target}/{category}/<name>.py; target-neutral abstractions stay at ir/{dialect}/{category}/. For example the whole MMA surface is target-owned — mma.py defines the TiledMma op, mma_atom.py defines MmaAtom / AtomPattern, and sm80_mma.py / wgmma.py define the CUDA instruction declarations. All four live under ir/tir/cuda/nn/. A target-specific rule for a target-neutral Op nests the same way: the CUDA MatMul type rule lives in ir/hir/cuda/nn/matmul.py beside no Op class of its own. The backend-bound construction stays in TIR: HIR is the checking reference side, and carrying the instruction name in that reference would make two GPU targets require different HIR references. (codegen/ and runtime/ are target-first instead — their primary axis is the target — so each tree is organized by its own primary axis.)

Rule 1d — a single carrier may be its category. When an IR category consists of one carrier operation named for that category, it may live directly at ir/<hir|tir>/<category>.py instead of the redundant ir/<hir|tir>/<category>/<category>.py. It MUST declare its dialect and category explicitly. ScheduleOp is this exception: it lives in ir/hir/schedule.py and registers as dialect tf, category schedule.

Rule 2 — one (node, target) codegen = one file. Each handler lives at codegen/<target>/tir/<category>/<name>.py. Stmt emitters, Expr-Op emitters, and tag-dispatched (arith, reduce) emitters each get their own file. HIR skeleton emission is the sole exception to the TIR input rule: its handler lives at codegen/<target>/hir/function.py and MUST NOT emit HIR operation bodies.

Rule 3 — what an IR-class file contains:

• HIR Op file (ir/hir/<cat>/<name>.py): Op class + @register_typeinfer(Op) + @register_cost_evaluator(Op) (if any).
• TIR effect Op file (ir/tir/<cat>/<name>.py): Op class + @register_typeinfer(Op) (returning UnitType) + @register_verify_stmt(Op). The verify rule keys on the Op class even though the invocation is an Evaluate(op, args) Stmt — see visitor-registry §5.
• TIR-owned Expr Op file (ir/tir/memory/{alloc_tensor,ptr_of,memory_span,tensor_view}.py, …): Op class + @register_typeinfer(Op) + @register_cost_evaluator(Op) (if any). Call-position constraints are checked by the enclosing Stmt's @register_verify_stmt.
• <category>/aliases.py file (Rule 1a): @register_alias(...) declarations whose builders construct the target Op instance.

Rule 4 — what a target codegen file contains: the register_codegen declaration naming that target, Role.EMIT, and that op / stmt class (visitor-registry §6), and nothing else. A target's own calling convention — how it declares a parameter and how a caller passes one — lives in that target's abi.py, not in a per-node file.

Rule 5 — <category>/__init__.py re-export rules: real Op submodules are re-exported; aliases are imported only for registration side effects; user imports go through parser §2.

Rule 6 — one pass = one file. A pass class lives in passes/transforms/<pass_name>.py; internal visitors / mutators stay in that file.

Rule 7 — what template files contain. codegen/<target>/templates/*.j2 carry boilerplate assembly only; emitters live in Python walkers.

2.1 Package export rule

tilefoundry.ir.types re-exports type classes and make_* constructors. Other functions are imported from the module that owns them; for example, make_mesh is available at the package surface, while separate is imported from tilefoundry.ir.types.mesh.

3. Multi-agent parallelism guarantee

The lock granularity is a single (node, target) pair. The naming rules in §2 imply that two agents working on different (node, target) pairs touch disjoint files; cross-cutting changes confine themselves to the owning directory.

4. DSL package layout

The author-facing surface is delivered as a namespace package. The two sub-packages (tf and T) use the OpSchema registry for their corresponding dialect and return an Op class or alias builder; unknown names raise AttributeError.

4.1 Built-in op-class location convention

For @register_op to auto-derive dialect + category, an Op class MUST live under src/tilefoundry/ir/<hir|tir>/<category>/<file>.py, with its module name matching that path. Outside it the decorator requires explicit values; the standalone ir/hir/schedule.py carrier therefore declares dialect="tf" and category="schedule".

4.2 .pyi stub regeneration

The .pyi stubs reflect registered schemas only. After adding a new @register_op / @register_alias, regenerate them with python -m tilefoundry.dsl regen.