refactor(ir): add AtomSched operation declarations

docs/spec/code-organization.md

@@ -18,10 +18,11 @@ truth for the directory's structure and invariants.
 | Directory | Owning spec | Contents |
 |---|---|---|
 | `ir/core/` | [core-ir](./core-ir.md) | Shared node algebra: `Module` / `Expr` / `Var` / `Constant` / `Tuple` / `Op` / `Call` / `Stmt` (base class) / `OpSchema` / `ParamDef` / call-graph and ownership queries / typed metadata attach-detach and diagnostics / `@register_op` / `@register_alias` / `op_registry` / `errors`. |
+| `ir/pattern/` | [core-ir](./core-ir.md) | 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](./types.md) | 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,shard_layout,mesh}.py` | [shard](./shard.md) | `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 and the algebra over them each in their own module. |
-| `ir/mesh_scope.py` | [shard](./shard.md) | Which scope a statement stands inside and what it admits: `merge_mesh`, `device_layout`, `covered_by_scope`, `check_topology`. Neither a type nor a visitor, so it sits beside `ir/isl_interop.py` rather than in either. |
-| `ir/constraints/` | [parser](./parser.md) | 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/types/{int_tuple,stride,layout,layout_algebra,shard_layout,mesh}.py` | [shard](./shard.md) | `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 and the algebra over them each in their own module; mesh construction, separation, and topology-bound checking stay with `Mesh`. |
+| `ir/mesh_scope.py` | [shard](./shard.md) | 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](./parser.md) | 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](./visitor-mutator.md) | `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](./types.md) | 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.md) | HIR Op layer; one subdirectory per category (`math/` / `tensor/` / `nn/` / `shape/` / `sharding/`). One real Op per `.py` ([§2](#2-file-naming-and-content-rules) rule 1); surface-alias schemas have no per-name file and live in each category's `aliases.py` ([§2](#2-file-naming-and-content-rules) rule 5). |
@@ -101,9 +102,9 @@ physical directory layout reflects that boundary directly.
   contracts are distinct even though both are consumed across the codegen
   boundary.
 
-`ir/constraints/`, `visitor_registry/`, and `dump/` are cross-cutting packages;
-their stable responsibilities are owned by [parser](./parser.md),
-[visitor-registry](./visitor-registry.md), and [inspection](./inspection.md),
+`ir/pattern/`, `ir/clause/`, `visitor_registry/`, and `dump/` are cross-cutting packages;
+their stable responsibilities are owned by [core-ir](./core-ir.md),
+[parser](./parser.md), [visitor-registry](./visitor-registry.md), and [inspection](./inspection.md),
 respectively. Their internal file layout is not a per-Op contract.
 
 ## 2. File naming and content rules
@@ -181,6 +182,13 @@ go through [parser §2](./parser.md#2-syntax-and-rules).
 **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

docs/spec/codegen.md

@@ -144,6 +144,8 @@ class CudaCodegenContext(CodegenContext):
     def bind_extents(self, params) -> None: ...
     def reset_barrier_ids(self) -> None: ...
     def alloc_barrier_id(self) -> int: ...
+    def reset_smem_base(self) -> None: ...
+    def smem_base(self) -> str: ...
     def dtype_to_cpp(self, dtype_name: str) -> str: ...
 
 
@@ -181,12 +183,30 @@ class CpuCodegenContext(CodegenContext):
     types alone. `launches` is the geometry each device function is called at,
     keyed by `id(fn)`, settled where the `Launch` was written
     ([passes §7.3](./passes.md#73-insert_default_host_entry)).
+  - `smem_base` declares and returns one byte-addressed dynamic shared-memory
+    base per kernel. A numeric `TensorView` address adds its byte offset before
+    converting the result to a CuTe shared-memory pointer.
   - A target subclass owns the type strings and hardware counters only it can
     state; a handler MUST reach them through the context rather than reading
     the IR for them. Other helpers MAY be added per target.
 
 ### 2.4 Effect Op dispatch
 
+CUDA value emission lowers `TupleGetItem` according to its index form. A constant
+index names the selected tuple element directly. A dynamic index requires the
+homogeneous tuple established by type inference and emits
+`cute::array{a, b, c}[index]`; dimension arithmetic in the index remains runtime
+C++ arithmetic.
+
+An IR `Tuple` is structural and has no target-side storage. A `LetStmt` binding
+one emits no C++ variable and continues with its body; its authored name (for
+example `lhs_stages`) therefore does not appear in generated C++. Each dynamic
+`TupleGetItem` use constructs its own `cute::array`, so indexing the same tuple
+N times constructs N arrays. Constant selection remains valid for heterogeneous
+tuples because it does not materialize an aggregate. The codegen context records
+the structural tuple by its fresh SSA `Var` identity so consumers can recover
+its elements without target-side storage.
+
 Effect Ops (`Copy`, `Fill`, `Mma`, `tir.nn.*`, ...) appear in Stmt
 position as `Evaluate(op, args)` rather than as Stmt subclasses. The
 walker matches `Evaluate` and dispatches on `type(callable)` through
@@ -440,4 +460,3 @@ variant runs is decided on the device.
     the call already carries -- the parameter the open axis expands into. A
     shape outside every variant's range is a call-contract violation and the
     kernel traps.
-

docs/spec/core-ir.md

@@ -482,6 +482,11 @@ class Op:
     (see [tir §2.3](./tir.md#23-tir-ops)).
 
 ```python
+class MemoryEffect(Flag):
+    READ = auto()
+    WRITE = auto()
+
+
 class ParamDef:
     """Declare one Op input or attribute.
 
@@ -491,17 +496,22 @@ class ParamDef:
         pattern: attribute; Optional input-type predicate.
         optional: attribute; Whether None is accepted.
         default: attribute; Call-site default or the required-value sentinel.
+        effect: attribute; Declared storage effect, or None when undeclared.
     """
 
     kind: Literal["input", "attribute"]
     annotation: type = field(default=object)
     pattern: Pattern | None = None
     optional: bool = False
     default: Any = MISSING
+    effect: MemoryEffect | None = None
 ```
 
 - constraints:
   - a single Op parameter descriptor; the order of input-kind ParamDefs fixes `Call.args` position.
+  - `MemoryEffect` is a `Flag` with `READ` and `WRITE`; an input may declare
+    either or both. `None` means undeclared, while the zero flag is invalid.
+    Attributes cannot carry a memory effect.
 
 Example:
 
@@ -592,62 +602,85 @@ and specialization dispatch.
 class Pattern:
     """Carry a reusable dispatch predicate."""
 
-    def match(self, subject) -> bool: ...
+    def match(self, subject, captures=None) -> Match | None: ...
 ```
 
 - constraints:
   - shared by parser dispatch (`ParamDef.pattern`) and specialization dispatch
     (`Function.specializations` / `PrimFunction.specializations`).
+  - a successful match returns a truthy `Match` carrying named captures;
+    failure is `None`.
+
+The implementation is split by responsibility under `ir/pattern/`:
+
+- `pattern.py` defines `Pattern` and the composable classes
+  `OrPattern`, `AndPattern`, `SequencePattern`, `CapturePattern`,
+  `ConstraintPattern`, `GuardPattern`, `SwitchPattern`, `RangePattern`,
+  `MultipleOfPattern`, `OneOfPattern`, `AttrPattern`, `BitsPattern`,
+  `LayoutPattern`, `SwizzlePattern`, `ComposedLayoutPattern`, `MeshPattern`,
+  `ShardLayoutPattern`, `ScalarPattern`, `TensorPattern`, and
+  `WildcardPattern`. It also owns the `Scalar` and `Tensor` singletons.
+- `match.py` owns matches, captures, symbolic resolution, layout-frame reading,
+  and the shared description helpers.
+- `constraint.py` owns cross-operand `Constraint`, `DistinctConstraint`,
+  `SameConstraint`, and `SameModesConstraint` values.
+- `utils.py` owns exact-layout construction plus specialization naming and
+  dimension lookup.
+
+`LayoutPattern` checks `forward` and `injective` over the whole flattened
+arrangement by default. With `per_mode=True`, it checks each top-level mode
+independently; `MeshPattern` requires this explicit form because each mesh
+level uses its own numbering space. `MeshPattern` never changes the supplied
+pattern implicitly.
 
 Two consumer surfaces:
 
 - **Parser dispatch** — `ParamDef.pattern` ([§2.3](#23-op)) is matched against an
   argument's `Expr.type` during overload resolution. Subclasses used:
-  `ScalarPat` (rank-0), `TensorPat(rank?, dtype?)` (non-scalar), and
-  `AndPat(parts)` (conjunction). Two singletons are exported as
-  convenience: `Scalar = ScalarPat()` and `Tensor = TensorPat()`.
+  `ScalarPattern` (rank-0), `TensorPattern(rank?, dtype?)` (non-scalar), and
+  `AndPattern(parts)` (conjunction). Two singletons are exported as
+  convenience: `Scalar = ScalarPattern()` and `Tensor = TensorPattern()`.
 - **Specialization dispatch** — patterns appearing in
   `hir.Function.specializations` ([hir.md §1.1](./hir.md#11-function))
   and `tir.PrimFunction.specializations` describe which runtime
   shape range a variant covers. The HIR→TIR lowering inspects each
   pattern's fields directly; it does not call `match`.
 
-### 3.1 `DimVarRangePat`
+### 3.1 `RangePattern`
 
 ```python
-class DimVarRangePat(Pattern):
-    """Match one sub-range of a named dimension.
+class RangePattern(Pattern):
+    """Match a closed integer range, optionally naming a dimension.
 
     Attributes:
-        dim_var: attribute; Name of the dimension.
-        lo: attribute; Inclusive lower bound.
-        hi: attribute; Exclusive upper bound.
+        dim_var: attribute; Specialization dimension name, or empty otherwise.
+        lo: attribute; Optional inclusive lower bound.
+        hi: attribute; Optional inclusive upper bound.
     """
 
     dim_var: str = ""
-    lo: int = 0
-    hi: int = 0
+    lo: int | None = None
+    hi: int | None = None
 ```
 
 - constraints:
-  - This is the per-variant sub-range for a named `DimVar`; `match(v)` is
-    `lo <= v <= hi` and ignores `dim_var`.
-  - `dim_var` MUST be a non-empty `str` — the name of the `DimVar` the
-    range applies to. The lowering resolves it to a runtime
+  - At least one of `lo` and `hi` MUST be stated. Each stated bound MUST be a
+    plain `int` (`bool` is rejected), and two stated bounds MUST satisfy
+    `lo <= hi`. `RangePattern(lo=k)` and `RangePattern(hi=k)` express the
+    one-sided relations formerly represented by separate pattern classes.
+  - A specialization states both bounds and a non-empty `dim_var`: the name
+    of the `DimVar` the range applies to. The lowering resolves it to a runtime
     `ShapeOf(param, axis)` by walking the enclosing function signature.
-  - `lo` and `hi` MUST be plain `int`s (`bool` is rejected).
-  - The interval is closed `[lo, hi]`; construction MUST satisfy `lo <= hi`. A single-point
-    range is `[k, k+1)`.
-  - `match(value)` returns `True` for an `int` value `v` iff
-    `lo <= v <= hi`. The `dim_var` field does not participate in
-    `match`.
+  - A two-sided interval is closed `[lo, hi]`; a single-point range is `[k, k]`.
+  - `match(value)` succeeds for an `int` value `v` iff every stated bound
+    admits it. The `dim_var` field does not participate in `match`.
   - The pattern references a `DimVar` by name only. The envelope of
     the named dim lives on the `DimVar(name, lo, hi)` itself (see
     [types.md §4](./types.md#4-dim--symbolic-shape-dimensions)); the
-    `DimVarRangePat` carries the per-variant sub-range. Envelope
+    `RangePattern` carries the per-variant sub-range. Envelope
     containment (`pattern ⊆ DimVar envelope`) is checked in
     signature context — by the `@tilefoundry.func` validator and the
-    HIR→TIR lowering — not by `DimVarRangePat.__post_init__`.
+    HIR→TIR lowering — not by `RangePattern.__post_init__`.
 
 ## 4. Shared operation kinds
 

docs/spec/hir.md

@@ -236,7 +236,7 @@ the Op typeinfer registry. `TypeInferVisitor` handles it directly as a
 - Within a `Function` signature, every occurrence of a same-name
   `DimVar` across `params` and `return_type` MUST agree on its
   `(lo, hi)` bounds; a disagreement is a verify error. A
-  `DimVarRangePat` specialization MUST anchor to a `DimVar` reachable
+  `RangePattern` specialization MUST anchor to a `DimVar` reachable
   from an input parameter and lie within that `DimVar`'s envelope
   (see **Shape dispatch and specializations** below).
 
@@ -290,15 +290,15 @@ freeze** below).
   `return_type`: a variant specializes the body, not the signature. A variant
   runs in the same execution domain as its base because both are owned by the
   same `Module`.
-- A variant carries exactly one `DimVarRangePat` in `specializations`.
+- A variant carries exactly one `RangePattern` in `specializations`.
   The canonical signature is
   `";".join(f"{p.dim_var}${p.lo}_{p.hi}" for p in specializations)`
-  (v0 allows only `DimVarRangePat`). Two variants of one base MUST have
+  (v0 allows only `RangePattern`). Two variants of one base MUST have
   distinct canonical signatures.
 
 *Envelope coverage.* A dispatched function's parameter
 `TensorType.shape` carries a `DimVar(name, lo, hi)` whose `(lo, hi)` is
-the dispatch envelope; `DimVarRangePat` references that `DimVar` by name.
+the dispatch envelope; `RangePattern` references that `DimVar` by name.
 The variants' closed ranges MUST **partition** the envelope — pairwise
 **disjoint** and jointly **complete** (their union is exactly the DimVar's
 half-open `[lo, hi)` envelope). Adjacent closed ranges are written
@@ -310,7 +310,7 @@ typeinferred, lowered, or evaluated as a body. Only its variants carry
 executable bodies. There is no base body to fall back to.
 
 *Dispatch resolution.* A `Call` whose target is a dispatch prototype
-(`variants != ()`) is a dispatch call: the variant whose `DimVarRangePat`
+(`variants != ()`) is a dispatch call: the variant whose `RangePattern`
 matches is selected and is the call's result. Evaluation selects from the
 call's concrete argument shapes; specialization selects from the caller's
 stated dimension bindings. Both use the same variant table. A shape outside
@@ -943,20 +943,25 @@ class Stack(Op):
 
 ```python
 class TupleGetItem(Op):
-    """Extract one field from a tuple-typed expression.
+    """Extract one field from a tuple-typed expression by scalar index.
 
     Attributes:
         tuple_value: input; Tuple-typed expression.
-        index: attribute; Static field index.
+        index: input; Scalar field index.
     """
 
     tuple_value: Expr
-    index: int
+    index: Expr
 ```
 
 - constraints:
-  - `tuple_value.type` MUST be `TupleType` and `index` MUST be in range.
-  - The result type MUST be exactly the selected field type.
+  - `tuple_value.type` MUST be `TupleType`. A constant `index` MUST be in range,
+    and the result type is exactly the selected field type.
+  - A dynamic `index` requires a non-empty tuple whose field types are all equal;
+    the result has that common field type.
+  - Access through a constant index reads the selected field's leaf span. Dynamic
+    access conservatively reads every leaf because the selected field is known
+    only at runtime.
 
 ##### Reshape
 ```python

docs/spec/inspection.md

@@ -256,12 +256,12 @@ variant as an `@<name>.specialize(pattern)` block in declared order:
 def f(x: Tensor[(S,), "f32"]) -> Tensor[(S,), "f32"]:
     pass
 
-@f.specialize(DimVarRangePat("S", 1, 3))
+@f.specialize(RangePattern("S", 1, 3))
 def small_sequence(x: Tensor[(S,), "f32"]) -> Tensor[(S,), "f32"]:
     ...
 ```
 
-The pattern prints in its constructor form (`DimVarRangePat("S", 1, 3)`;
+The pattern prints in its constructor form (`RangePattern("S", 1, 3)`;
 other `Pattern` subclasses fall back to `repr(pattern)`). The emitted binding
 mirrors the authoring surface ([parser.md §2.1](./parser.md#21-syntax));
 when an IR variant has no display label, the printer synthesizes a valid binding

docs/spec/parser.md

@@ -48,10 +48,10 @@ physical source-file coordinates with a one-based start column.
 
 ### 1.3 Tuple Binding Metadata
 
-For `a, b = producer(...)`, detached `TupleGetItem(index=0)` and
-`TupleGetItem(index=1)` lexical values carry the respective target Name spans (`a` and `b`) and
-matching `BindingMetadata`; later reads do not replace that identity. A multi-carry loop's
-derived projections carry the `for` statement span and their carry binding name.
+For `a, b = producer(...)`, detached `TupleGetItem` values with scalar inputs `0` and `1`
+carry the respective target Name spans (`a` and `b`) and matching `BindingMetadata`; later
+reads do not replace that identity. A multi-carry loop's derived projections carry the `for`
+statement span and their carry binding name.
 
 ### 1.4 Context and Diagnostics
 
@@ -83,6 +83,11 @@ lowered from, so a bound naming a mesh coordinate reads back as it was printed.
 
 ## 2. Syntax and Rules
 
+Tuple subscripting lowers `stages[index]` to `TupleGetItem(stages, index)`.
+Literal negative indices are normalized against the tuple arity before lowering.
+A non-literal scalar index is accepted only when type inference can establish one
+common field type for every possible result.
+
 ### 2.1 Syntax
 
 <!-- parser-grammar:start -->

docs/spec/runtime.md

@@ -1259,6 +1259,9 @@ helpers genuinely capable of host compilation (e.g. `local_tensor`,
 <!-- generated: ops-copy -->
 ```cpp
 // include/tilefoundry/runtime/cuda/ops/copy.cuh
+template <class TPointer, class TLayout>
+CUTE_HOST_DEVICE auto tensor_view(TPointer pointer, TLayout layout);
+
 template <class TSrc, class TDst>
 __device__ void copy(TSrc const &src, TDst &dst);
 
@@ -1267,6 +1270,11 @@ __device__ void copy_async(TSrc const &src, TDst &dst);
 ```
 <!-- /generated -->
 
+**`tensor_view`.** Rebuilds the CuTe tensor engine that a TIR `TensorView`
+describes from its typed pointer and emitted layout. Residency stays on the
+pointer engine; a later `make_shard_tensor` adds the logical global and shard
+layouts without recovering the source tensor.
+
 **`copy`.**
 
 One entry, taking tensors, with the transfer shape, the strides, the element