Heterogeneous targets · P2-2: execution model below spawn (grid/semantics, then roles)

P2-2 — Execution model below spawn (grid, then roles)

spawn on is a placement; for is sequential. No grid, block, warp, lane, or subcore, and no iteration-space annotation. Peak Blackwell throughput requires warp specialization (producer warps issuing TMA, consumer warps issuing tcgen05.mma, rendezvousing on mbarriers) — you cannot write a Blackwell GEMM in a language with no warps, and unlike the TPU you cannot delegate it (no XLA doing it for you). Largest single piece; most safely deferred.

Stage 1 — annotated grid (worth doing alone; unlocks pipelining on both machines)

spawn on(Topology::TensorCore) grid([n_batch, n_heads, n_kv])
  semantics([parallel, parallel, arbitrary])
{
  let b = grid_index(0);
  let h = grid_index(1);
  ...
}

arbitrary on the KV axis tells the compiler the online-softmax accumulation is order-dependent while batch/head are free — exactly Pallas's DimensionSemantics, and directly usable as gpu.launch_func grid dims.

Stage 2 — roles (sketch, B200-specific, genuinely open design)

spawn on(Topology::TensorCore5) grid([m, n]) semantics([parallel, parallel]) {
  role producer(warps: 1) {
    let t = transfer_async(tile(a, m, n), Memory::SMEM);   // P0-2 token
    yield_to(consumer, t);
  }
  role consumer(warps: 4) {
    let s = await_from(producer);
    mma_acc(s, Memory::TMEM);
  }
}

Stage 2 should not be committed to before P2-1 gives it a backend to land on; recorded so the stage-1 grid syntax is chosen with it in mind.

Where it lands

src/parser/stmt.rs (spawn grammar), src/dialect/VxDialect.cpp (grid/semantics attributes on vx.spawn), VxLowering.cpp (grid → gpu.launch_func dims). Depends on P0-2 (tokens) and P2-1 (backend).

From the Heterogeneous Target Gap Analysis (§9). Status tracked there and in the P0 decision log. Sibling to the landed P0 work (P0-1 fa056b4e, P0-3 33ce0b64, P0-4 5e10637a, P1-4a-core 34b31902).

Read docs/discussions/heterogeneous_target_gap_analysis.md §9 and docs/discussions/heterogeneous_p0_decision_log.md, then inspect src/parser/stmt.rs, src/dialect/VxDialect.cpp, and VxLowering.cpp. Resolve the P0-2 and P2-1 dependencies before evaluating the Stage 1 grid syntax. Done means annotated grid and semantics reach the dialect and lower to gpu.launch_func dimensions; Stage 2 roles remain an explicitly deferred design.