vt::RmsNorm: the kF16 ACTIVATION the seam admits and the CPU arm serves is refused by both device arms

Row: MODEL-MM-QWEN4-EXP

Spec: rmsnorm-gamma-dtype-twins.md, under ## Owed.
Found while closing #2503. Same defect class, one operand over, so it is not fixed in that flow.

The divergence

vt::RmsNorm admits IsFloat(x.dtype) (src/vt/ops.cpp:1030), and IsFloat is
{kF32, kF16, kBF16} (src/vt/ops.cpp:24). The CPU kernel widens a kF16 ACTIVATION
like any other -- WidenRowToF32(x.dtype, ...) at src/vt/cpu/cpu_ops.cpp:576, and the
explicit refusal above it at :561-562 lists f32/f16/bf16 as accepted. Both device arms
refuse it by name:

• src/vt/cuda/cuda_ops.cu RmsNormKernelCuda --- "cuda rmsnorm: unsupported input dtype (f32/bf16 only)"
• src/vt/rocm/rocm_rmsnorm.hip RmsNormKernelRocm --- "rocm rmsnorm: unsupported input dtype (f32/bf16 only)"

cuda_qwen4_exp.cu:60-62 states the rule this breaks:

a device arm that refused a dtype its CPU sibling accepts would be a divergence to record

It is the same shape as #2503 half 1, which was closed by SUPPORTING kF16 as a gamma.
The activation half was left open deliberately, because the two are not the same size of
change.

Why the gamma half was cheap and this one is not

The gamma is read once per element through Load(), which is overloaded per element type,
so a new gamma dtype is one switch case and one __device__ inline float Load(const __half*, ...). The ACTIVATION's type Tin also drives:

• ResRound<Tres> and the residual round-trip;
• LoadVec<N>, which has f32 and bf16 specialisations only;
• both vectorised decode-fast paths (TryLaunchRmsNormDecodeFast,
  TryLaunchRmsNormDecodeFastF32), whose guards are written against exactly two dtypes;
• the out.dtype switch, since IsOutFloat is {kF32, kBF16} and a f16 activation with
  a f16 output would need an output arm too.

Is it live?

Not today, as far as the loaders go: no model in this tree runs at a f16 model dtype on a
device backend. It is nevertheless a promise three surfaces make and one does not keep ---
CpuPlatform, CudaPlatform and RocmPlatform all return {kBF16, kF16, kF32} from
supported_dtypes() (src/vllm/platforms/cpu.cpp:22-24, cuda.cpp:111-113,
rocm.cpp:58-60), mirroring upstream interface.py:181-187, and
src/vllm/model_executor/models/qwen3_5.cpp:581 and
src/vllm/model_executor/models/nemotron_h_weights.cpp:941 both map a
"float16"/"half" config dtype to vt::DType::kF16. So a f16 checkpoint is
representable and would hit this refusal.

What would close this

Either instantiate the f32/bf16 kernel family for a f16 Tin and out --- including
guards on both fast paths so a f16 activation cannot enter a kernel written for two other
dtypes --- or narrow supported_dtypes() and the seam so the tree stops advertising a
model dtype no device kernel serves. The second is smaller and may be the honest one;
this issue does not pick.

Read .agents/specs/rmsnorm-gamma-dtype-twins.md under “Owed”, then compare the dtype checks and fast-path guards in src/vt/cuda/cuda_ops.cu, src/vt/rocm/rocm_rmsnorm.hip, and src/vt/cpu/cpu_ops.cpp. Check the platform supported_dtypes() declarations and the Qwen and Nemotron dtype mappings. Done means the advertised f16 path is served consistently by device arms, or the seam and supported dtype lists no longer advertise it.