ops: reject non-contiguous MoE weights

[morluto]

Fixes #3.

TL;DR

Reject non-contiguous MoE weight operands before they reach ThunderKittens' tensor_to_gl conversion and TMA-backed kernels.

Problem

The fused BF16 and MXFP8 wrappers checked weight shapes and devices but not contiguity. MoK's C++ entry points convert these operands directly with kittens::py::tensor_to_gl, bypassing the contiguity validation provided by ThunderKittens' higher-level object path.

Change

• Add one shared contiguity validator in mok/ops.py.
• Apply it to shared/routed weights and MXFP8 weight-scale tensors in all four fused forward/backward wrappers.
• Raise a named ValueError before the CUDA entry point is called; no implicit .contiguous() copy is introduced.

Review order

1. mok/ops.py: the validation boundary and the four affected operand sets.
2. tests/test_ops.py: one shape-preserving non-contiguous weight case per wrapper.

There are no generated or mechanical files in this change.

Regression coverage

The tests use as_strided to create a same-shaped non-contiguous shared weight and assert that each BF16/MXFP8 forward/backward wrapper raises ValueError during Python validation.

Validation

• python -m py_compile mok/ops.py tests/test_ops.py — passed
• git diff --check — passed
• Shape-preserving non-contiguous view fixture — passed
• Full distributed CUDA operator tests were not run in this environment because the compiled extension/GPU test runtime is unavailable.

---

Note

Low Risk
Defensive input checks only; fails fast on invalid layouts without changing kernel behavior for already-contiguous tensors.

Overview
Adds _validate_contiguous_tensors in mok/ops.py and calls it on shared/routed MoE weights (and MXFP8 scale tensors where applicable) in the four dispatch_mlp_swiglu_combine_* forward/backward wrappers, after shape checks and before _C / ThunderKittens tensor_to_gl paths.

Non-contiguous operands now raise a clear ValueError ("{name} must be contiguous"); there is no silent .contiguous() copy.

tests/test_ops.py adds an as_strided helper and one non-contiguous weight negative case per BF16/MXFP8 fwd/bwd wrapper.

^{Reviewed by Cursor Bugbot for commit 915468a. Bugbot is set up for automated code reviews on this repo. Configure here.}