claude
31d3ee99bb
fix: MoE kernel include paths — device_utils.cuh + arch_condition.h
...
Fixed xllm internal paths to our headers/ directory:
kernels/cuda/device_utils.cuh → device_utils.cuh
core/kernels/cuda/device_utils.cuh → device_utils.cuh
core/kernels/cuda/arch_condition.h → arch_condition.h (copied)
2026-08-14 11:31:54 +00:00
claude
df6a0f5d47
fix: remove cuda/functional from MoE topk kernels (not available on corex)
2026-08-14 11:29:19 +00:00
claude
a50adefdfc
feat: xllm MoE CUDA kernels — fused_topk + compute_index + combine
...
3 MoE kernel files adapted for corex:
moe_fused_topk.cu: LOG(FATAL)→TORCH_CHECK, +torch/extension.h
moe_compute_index.cu: CHECK_LE→TORCH_CHECK, uses cub::BlockScan (corex CUB)
moe_combine.cu: fixed duplicate include, +torch/extension.h
New pybind binding: xllm_moe_bind.cpp
→ moe_fused_topk(gating, topk, renormalize, bias, scoring_func)
→ moe_compute_index(expert_id, num_experts)
→ moe_combine_result(gemm2, weights, N, topk)
AST verification added for all 3 functions
2026-08-14 11:23:49 +00:00
claude
0359103b9b
fix: remove glog/cuda_ops_api.h dependency from all xllm CUDA kernels
...
cuda_ops_api.h includes glog/logging.h and ATen/DynamicLibrary.h
which are not available in corex standalone compilation.
All kernel .cu files only need device_utils.cuh (provides namespace,
XLLM_KERNEL_ATTR macro, CUB includes, type helpers).
Fixed files:
norm.cu, activation.cu, rope.cu, block_copy.cu, reshape_paged_cache.cu
moe/moe_combine.cu, moe/moe_compute_index.cu, moe/moe_fused_topk.cu
2026-08-14 10:30:10 +00:00
project6-dev
0b0c47fddd
fix(critical): fold max_completion_tokens + max_num_seqs=2 + max_model_len=80000 + xllm_latest layer import
...
Sub 655 root causes (confirmed from log analysis):
1. protocol.py: max_completion_tokens never folded into max_tokens
→ 162/881 replay requests rejected 400 (extra_forbidden)
2. max_num_seqs=1 → t2_n_2 test fails (needs n=2)
3. max_model_len=131072 → OOM crash at 62% replay, opencompass all 0
Fixes:
- protocol.py: model_validator fold_max_completion_tokens
- yaml: max_num_seqs=2, max_model_len=80000, PYTORCH_CUDA_ALLOC_CONF
- topk_softmax stays =0 (corex CUB BlockReduce incompatible on BI-V100)
xllm_latest import to ex_engine/:
- npu_torch layers: GDN(1164L), Qwen3.5 GDN, attention, fused_moe
- cuda/moe kernels: topk_softmax_kernels.cuh, moe_combine, moe_compute_index
- npu kernels: causal_conv1d, recurrent_gated_delta_rule
- model headers: qwen3_5.h, qwen3_next.h
2026-08-13 03:19:39 +00:00
