Module fused_row 

The fused row-op CUDA emitter — the real one-block-per-row shared-memory reduction kernel.

This generalizes the hand-written cuda/qwen_forward.cu::rmsnorm (one block per row, a strided partial into shared memory, a tree reduction, an rsqrt, a rescale) into a generated kernel parameterized by a recognized FusedRowSpec. It replaces the toy serial CUDA_RMSNORM_TEMPLATE / CUDA_SOFTMAX_TEMPLATE: those reduced on a single thread over one flat vector; this reduces per row in parallel over rows × cols, which is the transformer-normalization shape SOL-ExecBench scores. The flat sugar is the rows == 1 case.

It is chip-parametric: the provenance header records the target part (from CudaKnobs::chip), so the same recognized op emits a kernel stamped for the RTX 3060 or the B200. (Precision/bf16 loads and an extern "C" harness entrypoint layer on in the following steps; today the datapath is f32, matching the Real sugar it lowers, and the host prints Σ Y — the checksum contract the reduction lock reads.)

Functions

emit_fused_row

Emit a self-contained CUDA program for a recognized fused row op: one __global__ block per row doing a shared-memory tree reduction over cols, plus a host that fills X deterministically and prints Σ Y. knobs.chip stamps the provenance header (and drives the arch the caller compiles for).

emit_swiglu

Emit a self-contained CUDA program for a recognized SwiGLU map — y[i] = silu(gate[i]) * up[i], one thread per element, self-contained host (deterministic fills + Σ Y checksum). Chip-stamped like the fused-row.