36676f2d1b
SGEMM repos (upstream_ref/sgemm_cuda/, 41 files): siboehm/SGEMM_CUDA: kernel 1-12, runner, CMake, cuBLAS benchmark wangzyon/NVIDIA_SGEMM_PRACTICE: kernel 1-7 (Chinese comments), utils edtallison/sgemm-cuda: kernel 01-09 (learning notes), Makefile xllm kernels (ex_engine/xllm_kernels/cuda/): fused_qknorm_rope.cu + bind — saves 128 kernel launches/fwd xattention/ — 6 files from upstream xllm headers: corex_compat_utils.h, topk_last_dim.cuh ilu/CMakeLists.txt SO_BUILD_MANIFEST.md — complete .so inventory and call chain analysis
Upstream Reference: Deep-Spark xllm + vllm (FULL TREE)
Source repos (cloned 2026-08-09, Apache 2.0):
Deep-Spark/xllm— Iluvatar official C++ LLM inference engine (1470 files)Deep-Spark/vllm— Iluvatar official vllm fork (703 files, csrc + model layer)
What's here
xllm/ (complete source minus git/binaries/submodules)
天数智芯官方下一代推理引擎，C++ 原生，多平台(CUDA/ILU/MLU/NPU)。 包含 kernels → layers → models → runtime → scheduler → api_service 完整栈。
Key subtrees:
xllm/core/kernels/ilu/— ixformer API wrappers (ixformer.h是金矿)xllm/core/kernels/cuda/moe/— MoE CUDA kernels (topk_softmax, fused_topk)xllm/core/kernels/cuda/— activation, norm, rope, attention CUDA kernelsxllm/core/layers/ilu/— Iluvatar FusedMoE完整pipelinexllm/core/layers/npu_torch/— GatedDeltaNet C++ implementationxllm/models/llm/qwen3_5.h— Qwen3.5 model definitionxllm/compiler/tilelang/— GDN kernel code generation
ds_vllm/ (csrc + model layers + fused_moe)
天数智芯官方vllm fork，Python + CUDA torch extension。
csrc/— ALL CUDA source (attention, moe, quantization, cache)csrc/libtorch_stable/moe/topk_softmax_kernels.cu— vllm topk_softmaxvllm/_custom_ops.py— Python → torch.ops._moe_C bridgevllm/model_executor/models/qwen3_5.py— ds_vllm的qwen3_5实现vllm/model_executor/layers/fused_moe/— vllm FusedMoE Python layer