fix: apply x_stride1 in the rmsnorm normalize pass

[truong-v]

Fixes #1467

_rms_norm_fwd_fused honours x_stride1 when it accumulates the variance, and honours y_stride1 when it stores, but the normalize pass reads X + cols. For any input whose last-dim stride is not 1 the row's RMS is therefore computed from the right elements while the values it scales come from elsewhere, and the output is silently wrong.

-        x = tl.load(X + cols, mask=mask, other=0.0).to(tl.float32)
+        x = tl.load(X + cols * x_stride1, mask=mask, other=0.0).to(tl.float32)

rmsnorm_forward reaches the kernel with such a tensor without complaint: x.view(-1, x.shape[-1]) is shape-preserving for a 2-D input and succeeds whatever the strides are, and torch.empty_like(x) gives the output matching strides.

The gemma4, deepseek3_2 and NormWeight call sites all pass last-dim-contiguous tensors today, so no shipped model is affected — deepseek3_2's cache_kv[:, :, :kv_lora_rank] slice keeps stride 1 on the last dim. This makes the kernel consistent with the strides it already accepts.

Testing

New unit_tests/common/basemodel/triton_kernel/test_rmsnorm.py, 12 cases (3 shapes × with/without weight × contiguous/strided). All 12 pass with this change. Reverting the kernel and keeping the tests fails the 6 strided cases and passes the 6 contiguous ones.

Against torch_rms_norm from the same module, at M=64, N=256 float32:

input	before	after
contiguous, weight	0.0	0.0
contiguous, no weight	0.0	0.0
transposed (last-dim stride 64), weight	5.83 (15801/16384 cells wrong)	0.0
transposed (last-dim stride 64), no weight	6.58 (16230/16384 cells wrong)	0.0
the same values, made contiguous	0.0	0.0

Performance

triton.testing.do_bench on an NVIDIA B200, three interleaved before/after rounds:

shape	before	after
M=4096 N=4096 fp16	20.48 / 20.52 / 20.48 us	20.47 / 20.45 / 20.48 us
M=8192 N=8192 fp16	51.21 / 51.21 / 51.21 us	51.21 / 51.21 / 51.20 us
M=16384 N=2048 bf16	34.92 / 34.93 / 34.93 us	34.95 / 34.93 / 34.91 us
M=1024 N=7168 bf16	12.42 / 12.44 / 12.43 us	12.45 / 12.44 / 12.43 us

No measurable cost — the kernel is memory-bound and the variance pass already does the same multiply.

black --line-length=120 and flake8 with the repo's pre-commit arguments are clean on both files.

Environment

• ModelTC/lightllm at fe9bdabfc331b990124f1ec27daf6bb7945cf7ee
• NVIDIA B200 (sm_100), driver 595.71.05
• torch 2.13.0+cu130, triton 3.7.1, Python 3.12