spike: quantized-matmul / dequant-to-matmul fusion investigation on IREE (CUDA first)

[inureyes]

Part of #570

Summary

Investigate representations that let IREE fuse dequant into the matmul (or lower a quantized dot to a fused / int8 kernel) on CUDA, so the packed int8 path stops materializing a reconstructed f32 weight to DRAM every decode step. This spike de-risks Phase 2 by separating what is achievable in-tree from what is upstream IREE work, and picks the representation for #574.

Scope

Baseline: the #568 hand-rolled in-graph dequant (~678 decode dispatches, ~1.6 tok/s vs ~6.7 f16 on GB10). Investigate, and run real graphs through iree-compile + the runtime for, at least:

• MLIR quant dialect: !quant.uniform, stablehlo.uniform_dequantize.
• linalg.quantized_matmul.
• Data-tiling / ukernel paths.
• iree-compile flags beyond the three already tried (aggressive-fusion, generalize-matmul, early-trunc-fusion), which left the dispatch count at 678 -> 677.

For each representation, measure dispatch count + tok/s on GB10 CUDA.

Integration / acceptance

This is a spike, so the deliverable is a written finding rather than shipped default behavior, but it MUST run real graphs through iree-compile and the runtime (not paper analysis):

• A documented recommendation stating, per representation, the measured decode dispatch count and tok/s.
• A clear split of what fuses in-tree (authored in the portable graph / iree-compile flags) vs what needs upstream IREE work.
• The specific representation to implement in track(upstream-iree): fused dequant->matmul for the packed int8 bandwidth win #574, with the emitter / ABI changes it implies.

Dependencies

Depends on #571 (reproducible IREE toolchain) for the CUDA build and measurement.