Heterogeneous targets · R-2 (research): derive kernel parameters from the declared machine model

Part of hiraditya/Vx.1#256. Research direction — the strategic answer to "why not just use CuTe DSL?"

The asymmetry

CuTe knows layouts and instructions. It does not carry a machine model. In a CuTe kernel, "does this tile fit in SMEM" is verified by the author or discovered as a ptxas shared-memory-overflow error; capacities, bandwidths, granules and hop costs do not exist as data anywhere in that pipeline.

Vx has exactly that, declared and already statically checked (E6009/E6010, granule-rounded offsets, the transfer cost graph).

So the place Vx can plausibly exceed CuTe is not the instruction schedule — it is deriving the parameters a CuTe author hardcodes.

Candidates

1. Tile sizes. Given capacity: 228 KB, granule: 1 KB, scope: sm, the element type, and the working set (Q + K + V + S + O accumulator), the largest legal B_r/B_c is a constraint solve. The E6010 cumulative working-set check is the feasibility half of that problem already — this is inverting it from "does this fit" to "what is the largest that fits". CuTe authors pick 128×128 because someone benchmarked it.

2. Pipeline depth. stages = capacity / tile_bytes, both of which Vx knows. CuTe makes Stages a template parameter a human tunes.

3. Placement. HBM vs L2 vs SMEM vs TMEM given access counts and declared bandwidths is an optimisation over the transfer cost graph — i.e. the FA-4 paper's §3.1.1 roofline, which the paper performs informally in prose and Vx could compute.

4. Where the wait goes. The most interesting one: the seam checker is a latency-hiding optimiser in disguise. It currently proves "you must wait before this use." Inverting the same obligation yields the latest legal wait point — which is exactly what double-buffering does by hand. Depends on hiraditya/Vx.1#244 existing first.

The claim this supports

Parity on instructions, advantage on parameters. Both emit tcgen05.mma and cp.async.bulk.tensor; CuTe gets there with magic numbers a human tuned, Vx gets there by solving the constraint system it already type-checks.

That is a defensible answer to "why not just use CuTe", which the current hiraditya/Vx.1#256 positioning does not have.

Prerequisites and honest boundaries

• #246 (layout) is load-bearing. Every derivation above needs bytes-per-tile, which depends on layout and packing. Vx has device + memory; layout is the missing third, and this whole direction is blocked on it.
• A cost model beyond bandwidth is missing. Choosing between two legal schedules needs instruction latency/throughput, which nothing in Vx declares. Declared bandwidth gives the memory roofline; it does not say whether one MMA variant beats another.
• Do not try to out-derive the instruction schedule. Warp specialisation structure, tcgen05 variant selection, TMEM partitioning — the hardware knowledge in CUTLASS's collective mainloops is person-decades deep. Emit that structure; compete on the parameters around it.

Suggested first step

Take the smallest one that is fully determined by data Vx already has: pipeline depth from SMEM capacity and tile bytes (candidate 2), on the B200 model in the gap analysis. It needs no new solver, exercises whether the declared-machine-model numbers are actually accurate enough to drive a decision, and produces a number that can be checked against what a CUTLASS kernel picks for the same shapes.

If that lines up, candidates 1 and 4 are worth real investment. If it does not, the machine model needs more fidelity before any of this is trustworthy — which is itself the finding.

Start with the B200 machine model, the existing SMEM capacity data, and layout work in #246, since tile bytes are required for this derivation. Investigate whether pipeline depth can be computed from capacity and tile size without adding a new solver. Done means producing a checkable depth and comparing it with the corresponding CUTLASS choice.