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 |
|