Replace the if/else mess in convert_tt_layer_to_vllm_kernel with a
deterministic spec table mirroring mapper.py; only branch left is dense
vs sparse layer. out_buffers is now required (fail-loud on missing key),
experts quantize in a single grouped-triton call instead of a per-expert
loop. Slot allocation moved into the NIXL broadcast init so the
converter never allocates, and expert_lead / konig / ImportError fallbacks
are gone.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>

Weight broadcast over NIXL/UCX now completes on the 12-node prod setup
(8 trainer x 8 GPU + 2 prefill + 2 decode). Key fixes stacked here:

- Per-rank NIC pin actually applied: drop the sbatch-level
  UCX_NET_DEVICES=all so trainer ranks keep their PIX-attached NIC.
  GPUs 4-7 have three PIX NICs, one of which (mlx5_8) is DOWN on every
  trainer node, and UCX=all was ending up on it for GPU 4 and hitting
  NIXL_ERR_REMOTE_DISCONNECT.
- Allocate slot dtype to match vLLM for layernorm affine params and
  the expert-routing bias, which vLLM stores in fp32.
- Gather rather than per-shard when a source tensor is below a 2 MiB
  threshold, cutting per-rank handle count ~60%.
- Chunked drain (flush_every=100) to bound UCX queue depth.
- Trainer dist.barrier() after rank 0 sees NIXL_READY so no rank
  RDMA-writes before all inference engines have acked /pause.
- Peer/tag context propagated through wait() for diagnostics.

PyTorch expandable_segments:True hands tensors out of cuMemCreate +
cuMemMap virtual ranges. ibv_reg_mr on such VA succeeds, but the mlx5
HCA's MMU walk at WRITE time fails with "Local protection" (syndrome
0x4) because nvidia_peermem cannot pin a VA that spans multiple
cuMemCreate handles. UCX closes the EP and NIXL surfaces it as
REMOTE_DISCONNECT.

Fix: carve out the NIXL-registered slot buffers into a dedicated
CUDAPluggableAllocator + MemPool that calls cudaMalloc / cudaFree
directly. Everything else in the trainer keeps using expandable
segments, preserving the fragmentation mitigation.

Also preload libcudart with RTLD_GLOBAL because TileLang's stub
runtime wins dlsym(RTLD_DEFAULT) otherwise and its self-check aborts
the process the moment we enter the MemPool context. Wrapped in
try/except since CDLL can fail on hosts without a real CUDA runtime.

Wire_bw on the 12-node disagg run stays at 7.1-7.6 GB/s — same as
the expandable_segments=False baseline.

Split the old GLM-specific _Spec into two model-agnostic dataclasses
in a new trainer/models/conversion_spec.py:

- QuantizationSpec(destination_dtype, scale_suffix="") owns the
  transformation. Empty scale_suffix means plain copy_cast; non-empty
  means FP8 block quantize (2D or 3D dispatched from src.ndim).
  QuantizationSpec.apply(src, out, scale_out) is the single entry
  point for every spec type, including bf16 / fp32 dtype casts.
- ConversionSpec(dst, sources, cat_dim, quantization) owns routing.
  scale_name and per_source_scale_key moved from private helpers to
  public methods.

Drops the ad-hoc vllm_fp32_srcs set, the duplicated 2D/3D quantize-fn
dispatch, and the hardcoded .weight/_weight_scale_inv suffix in
modeling_glm_moe_dsa.py. Spec table entries are explicit — each fp32
layernorm and FP8 projection inlines its QuantizationSpec.

Moved the libcudart.so RTLD_GLOBAL preload to the top of
sparse_mla_{fwd,bwd}.py, before the import tilelang that was winning
the dlsym race and aborting the process. Previously it lived in
classic_cuda_pool.py which was imported too late.

Added docs/nixl-weight-broadcast.md covering the whole workflow
including the new class design and the platform gotchas.

12-node prod run stays at 7.0-7.6 GB/s wire / 18-22 GB/s net after
the refactor.

HSDP:
- NIXLWeightBroadcast runs the full protocol only on replica 0 (dp_replicate
  rank 0). Non-primary replicas skip TransportPlan / agent / SPG entirely.
- SPG world size drops to (dp_shard × cp) + inference_ws; trainer rank in
  the SPG is the dp_shard_cp axis rank, not the global rank.
- Slots index by dp_shard_cp rank too (via _shard_rank_and_size) so
  ShardedSlot.remote_chunk_idx and GatheredSlot's round-robin fan-out
  align with the SPG membership.
- broadcast_weights adds a second dist.barrier() at the end so non-primary
  replicas exit in lockstep with the primary's push. Without it they'd
  race ahead into the next step while the primary is still draining.
- Config plumbs per-replica trainer_world_size via _infer_nixl_trainer_ws.

EP partition assertion:
- ExpertSlot.from_spec now verifies num_local × fsdp × ep == total_experts
  from the DTensor global shape — catches misconfigured meshes (e.g. EP
  axis not dividing cleanly into dp_shard_mod_ep) at slot construction
  instead of as a silent routing bug.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>