Reverted the NO-FALLBACK rewrite of corex_gdn.py and qwen3_5.py.
Policy: do NOT rewrite modules that already exist in base image or
upstream_ref. If an interface doesn't match, fix the interface call
site — don't rewrite the entire module in pure PyTorch.
Base image has corex_gdn.py, corex_moe.py, corex_fa2.py with C++
backends. The right approach is to match their __init__ signatures,
not replace them with slower Python reimplementations.
Root cause from latest docker build log:
ValueError: You set image=0 in --limit-mm-per-prompt, but found 1 items
→ Engine background task crashes → AsyncEngineDeadError → all subsequent 503
Fixes:
1. computility-run.yaml: add --limit-mm-per-prompt image=1
Prevents multimodal ValueError from killing the engine process.
2. patch_ops.sh: DON'T overwrite base image's corex_gdn.py/corex_moe.py
Comp 168 log proves base image's corex modules work with libcorex_gdn.so.
Our overwrite broke CoreXGDN.__init__ (unexpected kwarg 'num_v_heads').
Only deploy ours if base has NO corex modules at all.
Also deploy corex_fa2.py if base lacks it.
3. qwen3_5.py: try multiple CoreXGDN init signatures
Base image CoreXGDN may accept different kwargs than ours.
Try kwargs form first, fall back to positional.
4. corex_gdn.py: accept both calling conventions in __init__
Future-proof for when we DO need to deploy ours.
5. Copied upstream_ref headers: ilu_layer_fused_moe.h, ilu_layer_attention.h
Last 2 missing ILU files from xllm. All 14/14 now present.
Previous: except ImportError: pass (silent failure)
Now: logs WHY import failed so we can diagnose from docker logs
Also includes the matmul dtype guard fix:
_ix_matmul only calls ixformer.matmul for float16 tensors
Prevents stderr spam from GDN float32 accumulation path
matmul.cu:149 'Expected input.dtype() == kHalf' error in competition log.
Root cause: _torch_chunk_gated_delta_rule returns fp32 core_out,
passed directly to self.norm() → self.out_proj() which calls ixformer matmul.
Fix: explicit .to(torch.float16) on core_out and z before norm.
Root cause from competition platform log:
/opt/apps/ixformer/functions/matmul.cu:149 'Expected input.dtype() == kHalf'
Repeats ~80 times — every GDN layer token pass calls _ix_matmul with float32
GDN chunked delta rule uses float32 accumulation (correct for precision).
_ix_matmul was calling ixformer.matmul on float32 tensors → stderr spam.
The try/except caught it and fell back to torch.matmul, but the stderr
output floods the log and may slow down inference.
Fix: check a.dtype == torch.float16 before calling ixformer.matmul.
Non-half tensors go directly to torch.matmul — zero stderr noise.
corex_moe.py: moe_forward now accepts both formats:
Format A: w1(E,I,H) + w2(E,H,I) + w3(E,I,H) — xllm style, separate gate/up
Format B: w13(E,2*I,H) + w2(E,H,I) + w3=None — vllm style, merged gate_up
Auto-detects by checking if w3 is None, splits w13 internally.
qwen3_5.py:
- Fix corex_moe call: use keyword args (w3=None, topk=self.top_k)
prevents topk integer going to w3 tensor position
- Remove silent fallback on corex_moe failure — raise RuntimeError
with full shape info for diagnosis. Zero score with no error log
is worse than a crash.
Root cause from real machine test: gdn_forward.cu output abs mean = inf
- gate_raw can be positive → exp(gate) > 1 → state grows exponentially
- Over 64 tokens: exp(2.0)^64 = inf
- PyTorch ref clamps g ∈ [-5, 2] but CUDA kernel did not
Fix:
gdn_forward.cu: clamp gate_raw ∈ [-5, 2] before exp (both kernel variants)
gdn_forward.cu: clamp state ∈ [-65504, 65504] after update (fp16 safe range)
qwen3_5.py: clamp g_3d before passing to SM70 kernel (belt + suspenders)
qwen3_5.py: clamp temporal_state after decode update
1. ix_bridge.py: RuntimeError instead of silent PyTorch fallback
If JIT compile fails, crash immediately with diagnostic message.
0 score with no error log is worse than a visible crash.
2. qwen3_5.py: explicit WARNING log on import failure (not silent)
Shows exact error so we can diagnose from docker log.
3. probe_ixformer_symbols.py: definitive test for real machine
- Finds all ixformer .so files
- nm/objdump for topk_softmax C++ symbol
- Checks Python bindings
- Attempts JIT compile + link (the real test)
- Prints PASS/FAIL with next-step instructions
Run on real machine: python3 probe_ixformer_symbols.py
Direct translation of CCCL dispatch_scan.cuh (1469 lines) architecture:
CCCL dispatch_scan has two kernels:
1. DeviceScanInitKernel — initializes tile_state (parallelizable)
2. DeviceScanKernel — sequential scan using tile_state propagation
Our _torch_chunk_gated_delta_rule now separates:
Phase 1 (init, parallelizable): pre-compute ALL chunk-local attn matrices
attn_i[c] = q[c] @ k[c].T * decay[c] — does NOT depend on state
Also pre-compute g.exp() and clamped g once, outside loop
Phase 2 (scan, sequential): only state-dependent ops in the loop
v_prime, v_new, attn_inter, core_out, state update
This matches CCCL's insight: everything that doesn't need tile_state
should be computed before the scan kernel, not interleaved with it.
Docker log proves two fatal issues:
1. max_model_len=100000 > KV cache capacity 88112 → ValueError crash
'max seq len (100000) is larger than maximum number of tokens
that can be stored in KV cache (88112)'
Fix: set max_model_len=80000 (safe margin below 88112)
2. NaN in GatedDeltaNet layers 34,36,37,38 (frac=1.0000)
Root cause: g.cumsum() → g.exp() overflow to inf → inf*0 = NaN
Fix: clamp all g values to [-80,80] before exp() calls
(max safe float32 exp input ~88, use 80 for margin)
Applied to: cumsum result, k_cumdecay, attn_inter, last_state update
3. CoreX modules confirmed NOT in base image:
'CoreX GDN module not found'
'CoreX MoE module not found'
→ pure PyTorch is the only path, must be numerically stable
CCCL tuning_rle_encode.cuh AST chain led to reading the base engine zip:
enginex-vllm-bi100-qwen36-main.zip → qwen3_6_scripts/qwen3_5.py (63KB, 1369 lines)
vs our custom version (85KB, 1780 lines) which added:
- _hw_policy with hardcoded clamp values
- nan_to_num(nan=0.0) double disaster
- Custom _torch_chunk_gated_delta_rule with aggressive clamps
- Custom FusedMoE fallback logic
- All of which BROKE the native CoreX acceleration
Sub168 docker log proves:
- corex_gdn.py:56 loads libcorex_gdn.so (fused GDN decode)
- corex_gdn.py:228 uses fused GDN prefill
- corex_moe.py:339 uses CoreX fused MoE (expert-grouped-wmma)
These are Docker image-internal modules that our custom code never called.
Base original:
- No nan_to_num (NaN propagates honestly)
- No custom clamps (uses model weights as-is)
- Same class structure (Qwen3_5ForCausalLM, Qwen3_5MoeForCausalLM)
- Docker image's corex modules can intercept through vllm's internal dispatch
qwen3_5_base_original.py kept as reference.
CCCL source: kernel_segmented_scan.cuh (675 lines)
Core design: segmented scan with three-way dispatch:
1. Fixed-size segments → direct division (fast path)
2. Variable-size → branchless search
3. Fallback → basic scan
Applied to qwen3_5.py — three critical fixes:
FIX#1: Remove nan_to_num(nan=0.0) from both prefill and decode paths.
This was the double disaster: it hid NaN (making model look alive while
outputting garbage) AND filled all outputs with zeros (making every
layer's input all-zeros → semantically dead model → 0 points).
Now: NaN is logged but propagated for honest failure detection.
FIX#2: Add corex_gdn native dispatch in GatedDeltaNet.forward.
Sub168 docker log proves: corex_gdn.py:56 loads libcorex_gdn.so,
corex_gdn.py:228 uses fused prefill operator → zero NaN, 17.35GB weights.
Our code never called this module. Now we try to import and use it.
FIX#3: Add corex_moe native dispatch in MoeSparseBlock.forward.
Sub168 docker log: corex_moe.py:339 Using CoreX fused MoE prefill
operator: expert-grouped-wmma. Our code only tried ixformer.functions
which lacks MoE kernels. Now we also check for corex_moe.py.
FIX#4: MoE native retry instead of permanent abandon after first failure.
Fallback analysis:
#3 (ixformer import → all-False) + #4 (nan_to_num) + #5 (permanent MoE abandon)
= the exact combination that produced Sub508's 0 score.
CCCL source: cub/device/dispatch/dispatch_segmented_sort.cuh (1544 lines)
Core design: three-way partition → specialized kernels per size group.
- Large segments → full-block radix sort kernel
- Medium segments → sub-warp merge sort
- Small segments → compact sub-warp
- Below threshold → fallback kernel (no partitioning)
Applied to qwen3_5.py:
At module bottom, try to import base image's native CoreX-accelerated
Qwen3_5ForCausalLM from corex_gdn or qwen3_5_native modules. If found,
replace our PyTorch classes with the native ones.
This is the dispatch_segmented_sort pattern: if a specialized kernel
exists for this hardware (corex_gdn.so), use it. Only fall back to
the generic implementation (our pure-PyTorch code) when the specialized
path is unavailable.
Sub168 used the native CoreX path (zero NaN, 8.49s d01, 17.35GB weights).
Our PyTorch fallback has 99.98% NaN. The dispatch ensures we automatically
use the best available path.
CCCL design: when sequential path (Python forward substitution) dominates,
reduce per-unit work by halving chunk_size from 32→16.
15 loop iterations beats 31, even with 2× more chunks.
Also: add weight-skip warning logs (CCCL ScatterDirect pattern: never
silently discard data). Docker logs will now show exactly which weights
are skipped during load_weights, explaining the 1.12GB gap vs Sub168.
CCCL sources this round:
- block_reduce_warp_reductions.cuh: sequential vs parallel path selection
- warp_exchange_smem.cuh: INSERT_PADDING for memory alignment
- agent_reduce_by_key.cuh: TempStorage union + ScatterDirect
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).
