Sub509 docker logs show 99.98-100% NaN rate in ALL GatedDeltaNet layers.
nan_to_num(nan=0.0) replaces them with zeros — entire DeltaNet layers
produce zero output, crippling model quality. This is root cause of:
- d03_tool_call FAIL (model too impaired to output <tool_call> XML)
- d01 content[0] (no content output, only 1085 reasoning tokens)
- d04 content[0] (same: reasoning but no content)
- 11x slower than opponent (model generates excessive tokens)
Five-layer fix based on CCCL overflow_cast_t pattern:
1. Pre-cumsum clamp: g.clamp(-0.5, 0.5) before cumsum (was: no pre-clamp)
- Limits cumsum growth to ±32 for chunk_size=64
- Post-cumsum clamp tightened from ±20 to ±12
2. A_log clamp tightened: [-20,20] → [-5,5]
- exp(5) ≈ 148 vs exp(20) ≈ 4.9e8
- Prevents extreme decay rates that feed into g
3. Forward substitution per-row clamp: ±1e4
- _forward_sub_lower was the primary NaN amplifier
- Each x[i] = rhs[i] + A[i,:i]@x[:i] now clamped
4. Cross-chunk state clamp: ±1e4
- last_state *= g_last_exp can blow up across many chunks
- Both prefill and decode paths protected
5. Decode temporal_state in-place clamp: ±1e4
- ts_flat.clamp_() after baddbmm_ state update
Also adds SUB509_DEEP_DIAGNOSIS.md with full root cause analysis.
sync_handler.cuh entire design (140 lines):
Centralized synchronization resource manager for GPU kernels.
Two-phase lifecycle:
Phase 1 (host, constexpr): registerResource(numStages) + registerPhase()
Declares what resources are needed. No allocation yet.
Phase 2 (device, once): clusterInitSync()
Initializes all mbarriers in one pass. After this, no more registration.
Key properties:
- Non-copyable, non-movable (single source of truth)
- Fixed-size arrays (mMaxNumResources=10) — no dynamic allocation
- Destructor asserts mHasInitialized (catch forgotten init)
- Block-strided barrier init (all warps participate)
Translation to MoeSparseBlock:
Previous: hasattr() checks in forward hot path to lazy-init _use_native_moe
Now: Pre-declare _use_native_moe=None in __init__ (Phase 1: registration)
First forward resolves it via _hw_policy (Phase 2: initialization)
Subsequent forwards: None-check is faster than hasattr()
Also pre-declare _moe_out_buf fields to avoid attribute creation in forward.
CCCL source: cub/cub/detail/warpspeed/sync_handler.cuh
Maps to: qwen3_6_scripts/qwen3_5.py (Qwen3_5MoeSparseBlock)
Full translation of thrust/benchmarks/bench/all_of/basic.cu pattern:
thrust::all_of uses short-circuit evaluation — once a mismatch is found,
it stops scanning. The benchmark's MismatchAt parameter (0.01, 0.5, 1.0)
shows that early detection at position 1% saves reading the other 99%.
Translation to NaN checks in GatedDeltaNet prefill/decode:
OLD: torch.isnan(result).any() — full tensor scan always (creates bool
tensor of same size, then reduces). If NaN found, does ANOTHER full scan
for mean(), then ANOTHER for nan_to_num. = 3 full passes.
NEW: Sample first 64 + last 64 elements. If sample is clean, skip all
3 full passes (the common case after overflow_cast clamp fix).
If sample detects NaN, proceed with full nan_to_num.
For decode (num_seqs=1, hidden_dim=2560): out has 2560 elements.
Sample check: 128 elements = 5% of tensor.
For prefill (seq_len=18K, hidden_dim=2560): result has 46M elements.
Sample check: 128 elements = 0.0003% of tensor.
On the happy path (no NaN), this eliminates O(N) work per layer per step.
dispatch_copy_mdspan.cuh entire design:
1. Check is_exhaustive() + have_same_strides() (layout compatibility)
2. Fast path: if contiguous, use DeviceTransform (1D memcpy-like kernel)
3. Slow path: if non-contiguous, use DeviceFor::for_each_in_extents
Translation to MoE segment loop:
After sorting tokens by expert_id, tokens routed to the same expert
often have consecutive original indices. When they do, hidden_states
slice is zero-copy (view) vs fancy indexing (allocates new tensor).
Check: tok_ids_seg[-1] == tok_ids_seg[0] + n - 1 (contiguous range)
Fast: hidden_states[first:first+n] (zero-copy slice)
Slow: hidden_states[tok_ids_seg] (gather with copy)
CCCL source: cub/cub/device/dispatch/dispatch_copy_mdspan.cuh
Maps to: qwen3_6_scripts/qwen3_5.py (_pure_pytorch_experts)
Full translation of cub/block/block_scan.cuh BLOCK_SCAN_RAKING_MEMOIZE strategy
to _torch_chunk_gated_delta_rule cross-chunk scan loop:
CCCL RAKING_MEMOIZE: 'preserve upsweep segment values in registers while
performing warp-synchronous scan, allowing downsweep not to re-read from
shared memory.'
Translation: precompute g_exp_full, g_last_exp, g_diff_exp tensors outside
the sequential cross-chunk loop. Loop body now uses indexed lookups into
precomputed tensors instead of calling exp() 3 times per chunk iteration.
For seq_len=100K with chunk_size=64: 1562 chunks × 3 exp() = 4686 exp() calls
eliminated from the hot loop. Replaced with 3 bulk exp() + tensor indexing.
Memory tradeoff (same as RAKING_MEMOIZE's register pressure):
+3 tensors of shape (batch, heads, num_chunks, chunk_size) float32
= 3 × 1 × 6 × 1562 × 64 × 4B ≈ 7MB (negligible vs 16GB model weights)
Also inherits overflow_cast protection: g is clamped to [-20,20] before
exp(), so precomputed values stay in safe float32 range.
thrust basic_vector.cu: device→host copy is batched (D = H, one memcpy).
Our MoE segment loop did int() per iteration — N separate GPU→CPU syncs.
Fix: .tolist() does ONE sync for all segment boundaries.
Maps to: qwen3_6_scripts/qwen3_5.py (_pure_pytorch_experts prefill path)
CCCL block_scan_raking.cuh: parallel prefix scan over C elements using
GPU-native raking threads, not sequential host-driven loops.
Our _forward_sub_lower was a Python for-loop over chunk_size=64 rows,
each launching a separate matmul kernel. This is 64 sequential kernel
launches per DeltaNet layer per chunk.
Fix: Use torch.linalg.solve_triangular (cuBLAS trsm) which solves
the entire (I-A)@X=RHS system in ONE kernel launch. Falls back to
the Python loop if cuSOLVER is unavailable on BI-V100.
CCCL source: cub/cub/block/specializations/block_scan_raking.cuh
Maps to: qwen3_6_scripts/qwen3_5.py (_forward_sub_lower)
CCCL tuning_radix_sort.cuh teaches: when one kernel in a chain is unavailable,
replace ONLY that kernel while keeping downstream native ops alive.
Our MoE chain: topk_softmax → moe_align_block_size → invoke_fused_moe_kernel
BI-V100 ixformer lacks vllm_moe_topk_softmax, which killed the ENTIRE chain
and forced 100% PyTorch fallback (_pure_pytorch_experts: 256x F.linear loop).
Fix: Add try/except in topk_softmax with PyTorch fallback (softmax+topk).
Now the chain can proceed to native align+invoke kernels if they exist.
Also: dont permanently disable native path after first failure — retry once.
CCCL source: catch2_test_device_radix_sort_pairs.cu + tuning_radix_sort.cuh
Maps to: _custom_ops.py (topk_softmax) + qwen3_5.py (MoE forward)
CCCL overflow_cast.h pattern applied to qwen3_5.py:
- Prefill gate: A_log.float().clamp(-20,20).exp() prevents NaN cascade
- Decode gate: same clamp before exp (was unprotected, unlike prefill path)
- Decode g_t: clamp_(-20,20) before in-place exp_() (was raw exp_())
Docker logs show 99.98% NaN in GatedDeltaNet layers — these unprotected
exp() calls are the root cause.
CCCL checked_allocator.cuh pattern applied to model_runner.py:
- Wrap model forward in try/except torch.cuda.OutOfMemoryError
- On OOM: empty_cache + gc.collect + retry once
- Competitor Sub168 died permanently at layernorm x.float() OOM
during replay (docker log evidence). This recovery keeps server alive.
Source: cccl_upstream/libcudacxx/include/cuda/__numeric/overflow_cast.h
Source: cccl_upstream/c2h/include/c2h/checked_allocator.cuh
Docker log reveals: 'NaN in prefill GatedDeltaNet layer 0 (frac=0.9998)'
Every DeltaNet (linear attention) layer produces 99.98% NaN values.
nan_to_num replaces them with zeros, destroying model output quality.
This is the root cause of d10_thinking_disable_ctk gibberish output.
Root cause: g.cumsum(dim=-1) accumulates unbounded gate logits.
When fed to exp(), large values overflow to Inf, which propagates
as NaN through subsequent matmul and forward_sub operations.
Fix: Clamp cumulative gate logits to [-20, 20] before any exp().
Range keeps exp in [~2e-9, ~5e8] — safe for float32 accumulation.
Inspired by CCCL dispatch_reduce_deterministic.cuh: numerical
stability requires bounded intermediate values (RFA pattern).
Also in this log:
- FusedMoE: 'vllm_moe_topk_softmax' not in ixformer → PyTorch fallback
(expected, cannot fix without BI-V100 kernel rebuild)
- OOM at end of sub168: 31.72 GiB GPU with 30.86 GiB allocated
CCCL input: dispatch_reduce_deterministic.cuh RFA pattern,
tuning_batch_memcpy.cuh (small=128t×4buf, large=256t×32B)
BI-V100 base image does not have libcusolver.so at:
/opt/sw_home/local/cuda/lib64/libcusolver.so
torch.linalg.solve_triangular requires cuSOLVER which is missing.
Replace with row-by-row forward substitution using only basic
matmul and indexing ops (torch.zeros_like, matmul, indexing).
The linear_attention gated_delta_rule solves (I-A)@X=RHS where A
is strictly lower-triangular. Forward sub: x[0]=rhs[0],
x[i]=rhs[i]+A[i,:i]@x[:i]. Mathematically equivalent.
CCCL source read: cub/device/dispatch/dispatch_reduce_by_key.cuh
- DeviceReduceByKey sorts input by key, pads to tile boundary, then
one fused kernel processes all key-value segments in parallel.
- This is architecturally identical to base engine's fused_moe.py:
moe_align_block_size (sort+pad) → invoke_fused_moe_kernel (one launch).
Discovery: _custom_ops.py (line 776-806) confirms ixformer HAS native MoE:
- ixf_F.vllm_moe_topk_softmax
- ixf_F.vllm_moe_align_block_size
- ixf_F.vllm_invoke_fused_moe_kernel (takes only BLOCK_SIZE_M config)
Previous code assumed 'ixformer lacks MoE kernels' and used _pure_pytorch_experts
(Python for-loop over 256 experts). This may have been wrong or outdated.
Change: MoeSparseBlock.forward now tries self.experts (FusedMoE native) first.
If the native kernel fails on BI-V100, it catches the exception, logs a warning,
and permanently falls back to _pure_pytorch_experts for that instance.
Impact if native works: one fused CUDA kernel vs 256× F.linear calls = massive
decode speedup. Impact if native fails: same behavior as before (fallback).
Qwen3.6-35B-A3B has 256 experts × top_k=8. The baseline prefill MoE:
for eid in unique_eids: # up to 256 iterations
tokens = hidden_states[tok_ids] # SCATTERED gather
F.linear(tokens, w13[eid])
Problem: hidden_states[tok_ids] creates a non-contiguous gather for each expert.
With 16384 tokens × 256 experts, this is 256 scattered gathers per layer.
Optimization (CCCL segmented-sort pattern):
1. Flatten all token-expert pairs: (T×K,) assignments
2. Sort by expert ID: tokens for same expert become CONTIGUOUS
3. Each F.linear gets contiguous input → much better memory access
4. Activation (silu × up) computed in ONE fused op across all pairs
5. index_add_ scatter-back is one kernel call
Memory access improvement:
Before: 256 × hidden_states[random_indices] → scattered HBM reads
After: sorted_tokens[start:end] → sequential HBM reads per expert
The expert loop still exists (can't batch variable-size GEMMs with F.linear),
but each iteration reads contiguous memory instead of scattered indices.
