claude
50a249e0a3
Revert "feat: batched MoE expert GEMM — replaces Python for-loop"
...
This reverts commit 06d7713db6
2026-08-14 11:47:37 +00:00
claude
06d7713db6
feat: batched MoE expert GEMM — replaces Python for-loop
...
ixformer probe results:
✗ moe_w16a16_group_gemm NOT in ixformer .so
✗ CUTLASS grouped GEMM needs cuda/std (variadic function error on corex)
✓ ixformer_linear EXISTS (fused matmul)
✓ torch.mm works (uses corex cublas)
Solution: moe_batched_gemm.cu
- C++ loop over experts (eliminates Python overhead)
- torch::mm for GEMM (corex cublas, not F.linear Python)
- Fused silu_and_mul CUDA kernel (not PyTorch ops)
- Weighted scatter-add in C++
- Skips empty experts (no wasted compute)
Integration in qwen3_5.py:
_USE_XLLM_MOE_GEMM dispatches to moe_experts_forward()
Falls back to Python for-loop if not available
Build: bash qwen3_6_scripts/build_xllm_kernels.sh
2026-08-14 11:43:46 +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
093bfb380f
feat: pybind11 bindings for xllm CUDA kernels
...
norm.cu compiled successfully on BI-V100 (only warning: fp8 __host__ attr).
Failed at import because no PYBIND11_MODULE — now fixed.
New bindings/ directory with 4 binding files:
xllm_norm_bind.cpp → rms_norm, fused_add_rms_norm
xllm_activation_bind.cpp → silu_and_mul, gelu_and_mul, act_and_mul
xllm_rope_bind.cpp → rotary_embedding
xllm_cache_bind.cpp → reshape_paged_cache, block_copy
Build script updated: each .so = kernel .cu + binding .cpp
2026-08-14 10:50:13 +00:00
claude
51cb90b9ab
fix: adapt xllm norm.cu for corex CUB (CUDA 10.2)
...
Key change: replace CCCL 3.6 types with corex CUB equivalents
- cuda::std::plus<> → cub::Sum
- cuda::maximum<> → cub::Max
- Remove #include <cuda/std/functional>
Test results from real machine (3/4 passed):
✓ __shfl_down_sync works on ivcore10
✓ manual SMEM+shuffle block reduce works
✓ corex CUB cub::BlockReduce<float,256> compiles and runs correctly (32640)
✗ CCCL 3.6 variadic function issue — corex clang rejects device variadic
Confirmed: use /usr/local/corex/include/cub/ for all kernel code
cccl_upstream is reference only, NOT compilable on corex
Build script: bash qwen3_6_scripts/build_xllm_kernels.sh
2026-08-14 10:19:51 +00:00
