Based on a3c45d3b (last known working docker build):
- yaml: max-num-seqs=2 (fixes t2_n_2), TOPK_SOFTMAX=1 (use prebuilt .so)
- yaml: keep max-model-len=131072, gpu-mem=0.90 (prevents OOM)
- yaml: NO LD_PRELOAD (libcccl not built during docker build)
- xformers: revert to Q-tiling only (flash_attn caused OOM at profiling)
- .dockerignore: exclude all non-essential files from context
- remove libcccl_allocator.so from git tracking
What stays from recent work:
- 14 prebuilt .so (including corex_gdn_chunk_recurrent)
- qwen3_5.py with .float() fix and chunk_recurrent support
- All vendor_overrides and CCCL preload source (for future use)
install_prebuilt_corex.sh asserted exactly 13 artifacts, but SHA256SUMS
now contains 14 (corex_gdn_chunk_recurrent.so was added in 9a52f057).
This mismatch causes Docker build to fail at:
'expected 13 prebuilt CoreX artifacts, found 14'
The .so's torch_chunk_gated_delta_rule() only accepts positional args:
(Tensor, Tensor, Tensor, Tensor, Tensor, int, Optional[Tensor], bool, bool)
But Python calls it with keyword args:
(q, k, v, g, beta, initial_state=, output_final_state=, use_qk_l2norm_in_kernel=)
This causes 'incompatible function arguments' crash during profiling
(determine_num_available_blocks), killing the engine before it starts.
Fix: _HAS_COREX_GDN_CHUNK = False, forcing Python _torch_chunk_gated_delta_rule.
This is what a3c45d3b effectively did (its .so wasn't compiled), explaining
why a3c45d3b works but aa4b4992 crashes.
31 files had Windows line endings (\r\n) from merge commit. This causes
patch_ops.sh replace_once() to fail: anchor strings use \n but file
content has \r\n, so no match → patch fails → docker build fails.
Also added .gitattributes to force LF for all text files going forward.
Sub655 root cause: OpenAIBaseModel had extra='forbid', rejecting
max_completion_tokens and reasoning_effort as 'Extra inputs not permitted'.
180/881 replay requests returned HTTP 400 instead of being processed.
Fix: extra='allow'. The fold_max_completion_tokens validator already
converts max_completion_tokens→max_tokens correctly. Unknown fields
like reasoning_effort are now silently accepted instead of 400'd.
Also resolved yaml merge conflict (keep upstream 0.80 gpu-mem, no LD_PRELOAD).
flash_attn_varlen OOMs at 4096 tokens, Q-tiling also OOMs (K tensor too large).
During profiling (BI100_IN_STARTUP_PROFILE=1), return zeros immediately.
Profiling only measures memory footprint, not output correctness.
Restore: chunked_prefill=on, max_num_batched_tokens=4096.
flash_attn_varlen_func allocates O(n²) temp memory for 4096 dummy tokens
during profile_run, causing OOM at gpu_memory_utilization=0.80.
BI100_IN_STARTUP_PROFILE=1 env var is already set by
patch_worker_startup_profile_guard.py during the synthetic forward pass.
Real inference requests still use flash_attn_varlen (much faster).
These 14 lines were added after a3c45d3b (last confirmed working build).
The cccl_preload build and corex extension compile steps may cause
docker build failure on the competition platform even with || fallback.
Reverting to the exact patch_ops.sh from a3c45d3b.
.cu and .sh source files remain in the repo for future use.
commit 4c365b8c added 1106 CCCL device-level headers (294K lines) to
qwen3_6_scripts/cccl_preload/include/. These are NOT used by the allocator
preload (which only needs cub/util_allocator.cuh + 288 transitive deps)
and cannot compile on corex CUDA 10.2 anyway.
The extra headers doubled docker context from 15MB to 31MB, likely
causing platform build timeout or size limit failure.
Restoring to the original 288-header set that is proven to compile and run.
Two problems from real BI-V100 build:
1. 'CUDA versions below 12 are not supported'
→ Add CCCL_IGNORE_DEPRECATED_CUDA_BELOW_12 (official suppress macro)
2. corex thrust/complex.h conflicts with CCCL thrust headers
→ Split into two compilation units:
- cccl_moe_sort_scatter.cu: CCCL headers only, C API, no torch
- cccl_moe_sort_scatter_pybind.cpp: torch headers only, no CCCL
Same pattern as proven cccl_allocator_preload.cu
3. Variadic device functions rejected by corex clang:
→ is_referenceable.h: __test(...) → __test(long)
→ invoke.h: __any(...) → template __any(_T)
→ conjunction.h: __and_helper(...) → __and_helper(long)
SFINAE still works: int overload wins, long is fallback.
CCCL latest requires CUDA 12+, corex is 10.2. Device-level CUB headers
(DeviceRadixSort etc) pull in thrust/detail/type_traits.h which conflicts
with corex's thrust/complex.h namespace.
Rewrite to use block-level CUB BlockScan only (same pattern as the proven
corex_moe_index_combine.cu): histogram + prefix_sum + scatter.
No extra_include_paths needed — uses corex's built-in cub/block/block_scan.cuh.
CUB CachingDeviceAllocator::DeviceAllocate calls cudaMalloc internally
on cache miss. Without a guard, our intercepted cudaMalloc recurses
into DeviceAllocate → cudaMalloc → DeviceAllocate → segfault.
thread_local g_in_allocator flag detects reentrant calls and forwards
them directly to the real cudaMalloc/cudaFree via dlsym(RTLD_NEXT).
CUB_DISABLE_NAMESPACE_MAGIC requires CUB_WRAPPED_NAMESPACE.
CUB_WRAPPED_NAMESPACE=cccl_preload wraps cub into cccl_preload::cub.
Source must use cccl_preload::cub::CachingDeviceAllocator.
CUB_WRAPPED_NAMESPACE=cccl_preload wraps cub into cccl_preload::cub
but cccl_allocator_preload.cu uses bare cub:: — compilation fails.
_CCCL_COMPILER_GCC=1 conflicts with CCCL auto-detection (redefined warning).
Drop both. CUB_DISABLE_NAMESPACE_MAGIC alone is sufficient.
Route C: replace PyTorch's cudaMalloc/cudaFree with CCCL CUB's
CachingDeviceAllocator via LD_PRELOAD. Eliminates driver-level allocation
overhead by reusing freed GPU memory from a bin-based cache.
Based on cccl_upstream/cub/cub/util_allocator.cuh (901 lines).
Self-contained .so with no CCCL header dependencies at compile time.
Files:
- cccl_preload_allocator.cu: the allocator (405 lines)
- build_cccl_preload_allocator.sh: build script (corex clang++ or g++ fallback)
- test_cccl_preload.sh: smoke test suite for BI-V100
- patch_ops.sh: build during docker build
- computility-run.yaml: LD_PRELOAD env var for runtime
Config via env:
CCCL_ALLOC_BIN_GROWTH=8, MIN_BIN=3, MAX_BIN=13, MAX_CACHED_MB=4096
Test on real machine:
cd qwen3_6_scripts && bash test_cccl_preload.sh
- patch_ops.sh: build corex_gdn_chunk_recurrent.so alongside moe_index_combine
- qwen3_5.py: import corex_gdn_chunk_recurrent, use C++ version for prefill
chunks instead of Python _torch_chunk_gated_delta_rule
- C++ version from xllm upstream avoids Python loop overhead and has proper
fp32 accumulation (key for NaN prevention on BI-V100)
- Falls back to Python version if .so not available
Verified on real BI-V100:
flash_attn_func works with head_dim=256 (diff < 0.004, no NaN)
flash_attn_varlen_func works for variable-length batching
seq=1024: 1.7x faster than PyTorch matmul
The profiling-stage _run_sdpa_fallback now tries flash_attn_varlen_func
first, falls back to Python Q-tiling only on exception.
This addresses the 10-50x attention slowdown identified in the analysis:
Python Q-tiling: O(L^2) per-tile matmul in Python loop
flash_attn: fused kernel, O(L) memory, hardware-optimized
Extracted torch_chunk_gated_delta_rule and torch_recurrent_gated_delta_rule
from xllm_latest/core/layers/npu_torch/qwen3_gated_delta_net_base.cpp.
Pure PyTorch C++ — no NPU/ACL deps, no custom CUDA kernels.
Same algorithm as our Python _torch_chunk_gated_delta_rule but
avoids Python interpreter overhead in the chunk loop.
Verify on real BI-V100: python3 verify_gdn_cpp.py
Verified on real BI-V100:
moe_compute_index: 11.48x speedup (0.035ms vs 0.397ms)
moe_combine_result: 2.66x speedup (0.022ms vs 0.059ms)
Integration:
- qwen3_5.py: import corex_moe_index_combine, use in prefill path
with _USE_COREX_MOE_INDEX_COMBINE flag (env BI100_MOE_COREX_INDEX_COMBINE)
Falls back to PyTorch argsort+bincount if .so unavailable
- patch_ops.sh: compile corex_moe_index_combine.cu during docker build
