sglang-kernel（sgl-kernel）内核库完全指南：安装构建、新增算子开发流程与体积分析-北京尧图网络

sglang-kernelsgl-kernel内核库完全指南安装构建、新增算子开发流程与体积分析【免费下载链接】sglangSGLang is a high-performance serving framework for large language models and multimodal models.项目地址: https://gitcode.com/GitHub_Trending/sg/sglangsglang-kernel 是 SGLang 项目中面向 LLM 推理引擎的高性能计算内核库为大型语言模型与视觉语言模型提供经过深度优化的自定义算子kernel。本文以仓库内 python/sglang/kernels/aot/README.md 为主体结合源码、构建脚本与测试用例系统讲解其安装方式、从源码构建的方法、新增内核算子的完整开发流程、基准测试规范以及编译产物体积分析方法帮助开发者快速上手使用并参与内核贡献。sglang-kernel 是什么sglang-kernel此前名为sgl-kernel是 SGLang 项目提供的一套针对 LLM 推理引擎的优化计算原语库。它通过自定义 CUDA kernel 实现高效推理覆盖注意力MLA decode、FlashAttention、FlashMLA、GEMMFP8 / INT8 / AWQ / GPTQ、MoEtopk、FP8 blockwise、W4A8、归一化RMSNorm、fused add RMSNorm、旋转位置编码RoPE、采样、投机解码等大量算子。从源码结构看该库的完整源码树位于 python/sglang/kernels/aot 目录下而 Python 导入路径仍然保持为sgl_kernel这保证了与历史版本兼容import sgl_kernel from sgl_kernel import fp8_scaled_mm, cutlass_mla_decode, moe_align_block_size__init__.py见 python/sglang/kernels/aot/python/sgl_kernel/init.py集中导出了全部算子接口并针对不同平台做了分支处理在 macOS Apple Silicon 上仅导出 Metal 扩展不加载 CUDA 算子在 CUDA 环境下通过 load_utils.py 中的_load_architecture_specific_ops()根据 GPU 计算能力动态加载对应架构的编译产物在 ROCmtorch.version.hip、MUSAtorch.version.musa环境下分别补充对应的算子实现。值得一提的是load_utils.py 中的架构检测逻辑显示计算能力为 90即 Hopper/H100时加载sm90子目录下的common_ops采用 fast math 优化其他计算能力则加载sm100子目录下的版本采用 precise math 以保证兼容性。这与 CMakeLists.txt 中双库构建的设计一一对应详见下文「构建系统剖析」。安装版本要求安装sglang-kernel前需要确认本机环境满足PyTorch 2.13.0该版本为强绑定依赖见 pyproject.toml 中的dependencies [torch2.13.0]Python ≥ 3.10通过 pip 安装# 安装最新版本 pip3 install sglang-kernel --upgrade从 pyproject.toml 可以看到该包以sglang-kernel为发行名构建后端为scikit-build-core并会随包安装python/sgl_kernel下的 Python 源码与编译生成的.so扩展。仓库还针对不同硬件平台提供了独立的构建配置如 pyproject_rocm.toml、pyproject_cpu.toml、pyproject_musa.toml以及 setup_rocm.py、setup_musa.py、setup_metal.py 等平台适配脚本。从源码构建构建依赖从源码构建需要以下工具链CMake ≥ 3.31Python ≥ 3.10scikit-build-coreninja可选使用 Makefile 一键构建进入 python/sglang/kernels/aot 目录后直接执行make build从 Makefile 的build目标可以看到其内部完整流程先安装依赖scikit-build-core、isort、black再初始化并更新 git submodule随后通过uv build --wheel构建 wheel 包最后将产物安装到当前 Python 环境build: install-deps submodule ## Build and install wheel package rm -rf dist/* || true \ CMAKE_POLICY_VERSION_MINIMUM$(CMAKE_POLICY_VERSION_MINIMUM) \ MAX_JOBS$(MAX_JOBS) \ CMAKE_BUILD_PARALLEL_LEVEL$(CMAKE_BUILD_PARALLEL_LEVEL) \ CMAKE_ARGS$(CMAKE_ARGS) \ uv build --wheel -Cbuild-dir$(UV_BUILD_DIR) . --verbose --coloralways --no-build-isolation \ pip3 install dist/*whl --force-reinstall --no-deps限制构建资源占用CPU / 并行度默认情况下make build会使用本机全部 CPU 核心MAX_JOBS默认取nproc的值。在低内存主机上过高的并行度容易导致编译 OOM可以通过以下方式显式控制并行度# 限制并行任务数同时控制 make 与 cmake 的并行度 make build MAX_JOBS2 # 同时限制 NVCC 内部编译线程数进一步降低 CPU 占用与峰值内存 make build MAX_JOBS2 CMAKE_ARGS-DSGL_KERNEL_COMPILE_THREADS1关于SGL_KERNEL_COMPILE_THREADS的说明在 CMakeLists.txt 中该选项默认值为 32用于控制传给 NVCC 的--threadsN编译线程数对应SGL_KERNEL_CUDA_FLAGS中追加的--threads${SGL_KERNEL_COMPILE_THREADS}。CMake 注释明确提示此前直接写死--threads32会破坏CMAKE_BUILD_PARALLEL_LEVEL对并行度的控制并在低内存主机上触发 OOM因此才将其提取为独立选项。该选项要求必须是正整数否则构建会直接报错SGL_KERNEL_COMPILE_THREADS must be an integer。Makefile 其他常用目标Makefile 还提供了一组常用开发目标目标作用make install以开发模式安装pip install -e . --no-build-isolation需先初始化 submodulemake clean清理build、dist、*.egg-info构建产物make rebuild先 clean 再重新构建安装make test遍历tests目录逐个运行test_*.pymake format用 clang-format / isort / black / pre-commit 统一格式化源码make update new_version同步更新多个文件中的版本号make tree打印项目目录结构排除构建产物make ln生成compile_commands.json编译数据库构建系统剖析CMakeLists.txt 是理解整个构建流程的关键它揭示了几个值得关注的设计第三方依赖全部通过 FetchContent 拉取包括 NVIDIA cutlass、fmtlib、tritonv3.7.1、flashinfer、flash-attentionsgl-project/sgl-attn。每个依赖都固定了 commit 并校验 SHA256 哈希保证构建可复现。GITHUB_ARTIFACTORYCMake 选项允许将拉取地址替换为内部镜像例如内网 artifactory。按 CUDA 版本分支生成 gencode脚本检测CUDA_VERSION按 13.0 / 12.8 / 12.4 / 12.1 / 11.8 分级处理。默认启用 SM90 及以下架构compute_80/sm_80、compute_89/sm_89CUDA ≥ 12.8 时追加sm_100fCUDA ≥ 12.9或sm_100a与sm_120a在 aarch64 平台默认关闭 SM90 以下架构并追加 sm_87 / sm_110a / sm_121a 等。可选地SGL_KERNEL_ENABLE_SM90A、SGL_KERNEL_ENABLE_SM100A可以显式开启对应代际。双库构建SM90 与 SM100同一个SOURCES列表会编译出两个common_ops扩展——common_ops_sm90_build使用-use_fast_math并输出到sgl_kernel/sm90子目录common_ops_sm100_build使用精确数学并输出到sgl_kernel/sm100子目录。运行时由load_utils.py依据 GPU 计算能力选择加载哪一个。可选编译特性开关SGL_KERNEL_ENABLE_BF16默认 ON、SGL_KERNEL_ENABLE_FP8默认 ON含 E4M3/E5M2、SGL_KERNEL_ENABLE_FP4默认 OFF、SGL_KERNEL_ENABLE_FA3CUDA ≥ 12.4 且非 aarch64 时默认 ON、SGL_KERNEL_ENABLE_FA3_SPARSE_MASK默认 OFF。FA3 编译时还会根据ENABLE_BELOW_SM90等条件生成 sm80/sm86/sm90 的实例化源码。独立扩展模块除common_ops外还会构建flash_opsFA3、infllm_opsInfLLM-V2 FlashAttention仅编译 hdim 64/128 的 bf16 forward用于推理场景、spatial_opsgreen context 流并通过 cmake/flashmla.cmake 额外集成 FlashMLA同时把 triton 的 Python 内核目录一并安装到包内。新增一个 kernel 的完整流程六步开发流程README 给出了向sglang-kernel贡献新内核的标准流程对应仓库中的实际文件为实现 CUDA kernel在 csrc 目录下实现算子主体例如 csrc/gemm/fp8_gemm_kernel.cu、csrc/moe/moe_topk_softmax_kernels.cu。在头文件中暴露接口在 include/sgl_kernel_ops.h 中声明 C 函数原型该文件共 886 行按csrc/allreduce、csrc/attention、csrc/gemm、csrc/moe等来源分组组织。注册 torch extension在 csrc/common_extension.cc 中通过TORCH_LIBRARY_FRAGMENT注册算子。更新构建列表在 CMakeLists.txt 的SOURCES变量中追加新的 CUDA 源文件注释要求按字母序排序。暴露 Python 接口在 python/sgl_kernel 目录下新增 Python 模块并在init.py 中导出。添加测试与基准在 tests 添加 pytest 用例在 benchmark 添加 triton benchmark。使用 m.def 与 m.impl 注册算子创建 torch extension 时需要用m.def定义函数签名schema再用m.impl绑定设备实现。README 特别强调带完整 schema 的m.def是torch.compile正常工作的前提而不仅是简单的自动推导// 需要带 schema 的 def供 torch.compile 使用 m.def( bmm_fp8(Tensor A, Tensor B, Tensor! D, Tensor A_scale, Tensor B_scale, Tensor workspace_buffer, int cublas_handle) - ()); m.impl(bmm_fp8, torch::kCUDA, bmm_fp8);对照 common_extension.cc 中的实际写法可以看到同样的模式贯穿整个注册过程例如m.def( cutlass_mla_decode(Tensor! out, Tensor q_nope, Tensor q_pe, Tensor kv_c_and_k_pe_cache, Tensor seq_lens, Tensor page_table, Tensor! workspace, float sm_scale, int num_kv_splits) - ()); m.impl(cutlass_mla_decode, torch::kCUDA, cutlass_mla_decode);这里的Tensor!表示原地写入mutable参数Tensor?表示可空张量这些约定与 PyTorch 官方 schema 语法一致。用 make_pytorch_shim 适配 C 原生类型许多第三方 C 库的接口使用int、float这类原生类型而 PyTorch 绑定由于 Python 类型映射的原因要求使用int64_t和double。如果直接绑定会导致类型不匹配。为此仓库提供了 include/sgl_kernel_torch_shim.h 中的make_pytorch_shim工具函数自动完成转换。以int→int64_t为例通过特化模板声明转换规则并在转换时做范围校验// 添加 int - int64_t 的类型转换 template struct pytorch_library_compatible_typeint { using type int64_t; static int convert_from_type(int64_t arg) { TORCH_CHECK(arg std::numeric_limitsint::max(), value too large); TORCH_CHECK(arg std::numeric_limitsint::min(), value too small); return arg; } };然后直接包裹你的函数完成注册// 包装你的函数 m.impl(fwd, torch::kCUDA, make_pytorch_shim(mha_fwd));除了int/floatsgl_kernel_torch_shim.h 还处理了std::optionalT、std::optionalconst at::Tensor等 PyTorch 绑定不支持的类型形态这保证了第三方库如 flash_attn 的flash_api.cpp可以以最小改动接入便于后续与上游同步。测试与基准测试规范单元测试在 tests 目录下添加 pytest 测试。若某些测试需要按条件跳过例如新硬件特性尚未就绪必须使用pytest.mark.skipif并写明跳过原因pytest.mark.skipif( skip_condition, reasonNvfp4 Requires compute capability of 10 or above. )仓库中的真实示例可参考 tests/test_fp8_gemm.py它实现了torch_scaled_mm参考实现将输入升为 fp32 做 matmul 再乘 scale再与sgl_kernel.fp8_scaled_mm的结果逐项对比验证精度并通过_cuda_version_at_least等辅助函数判断运行环境是否支持原生标量 scale。此外 tests/conftest.py 提供共享的测试夹具与工具函数。性能基准在 benchmark 目录下使用 triton benchmark 框架编写基准脚本该目录已包含bench_fp8_gemm.py、bench_rmsnorm.py、bench_rotary_embedding.py、bench_cutlass_mla.py、bench_top_k_top_p_sampling.py等二十余个基准。README 明确建议推荐使用triton.testing.do_bench_cudagraph做内核基准测试。与triton.testing.do_bench相比do_bench_cudagraph的优势在于降低 CPU 开销对测量结果的影响内核性能数据更准确将 PDLProgrammatic Dependent Launch程序化依赖启动的影响纳入单个内核的结果中在支持 PDL 的架构SM ≥ 90上给出更贴近真实场景的性能数据。运行测试套件在 python/sglang/kernels/aot 目录下执行make test底层等价于对tests下所有test_*.py逐个执行python3。内核体积分析排查模板实例化膨胀编译产物wheel中的 CUDA kernel 尺寸分析用于识别过大的内核和模板实例化膨胀问题。该工具依赖cubloaty# 安装 cubloaty pip install cubloaty # 分析一个 wheel 文件 python analyze_whl_kernel_sizes.py path/to/sglang_kernel-*.whl # 自定义输出文件 python analyze_whl_kernel_sizes.py path/to/sglang_kernel-*.whl --output my_analysis.txt从 analyze_whl_kernel_sizes.py 的实现可以看到其工作原理解压 wheel → 递归查找所有.so与.cubin文件 → 对每个二进制文件调用cubloaty --format json→ 汇总所有 kernel 的名字与字节数。它会生成一份文本报告包含汇总统计kernel 总数、总体积、平均大小按名称前缀分组的 kernel 组Top 20按体积降序用于定位哪类算子的实例化最占空间单个 kernel 明细Top 20按体积排序列出所在.so文件、kernel 名称与体积占比。extract_kernel_prefix会把形如kernel_name模板参数...的名字截断到之前从而把同一模板的不同实例归并到同一组方便一眼看出模板展开带来的体积膨胀。使用建议在发布新版本 wheel 前后各跑一次体积分析观察新增 kernel 或模板特化对包体积的影响及时发现异常膨胀。常见问题FAQQCUDA 12.6 下出现 Segmentation fault段错误怎么办A将 ptxas 更新到 12.8 版本。该问题属于 CUDA 12.6 工具链的已知兼容性问题与 FlashAttention 相关栈在旧版 ptxas 下的代码生成缺陷有关升级 ptxas 即可规避。总结sglang-kernel是 SGLang 推理引擎的性能底座通过预编译的 CUDA 内核为注意力、GEMM、MoE、采样等关键路径提供优化原语。本文从安装、源码构建、资源控制、算子开发流程、类型适配、测试基准到体积分析完整梳理了围绕该库的工程实践要点。对于想为 SGLang 贡献内核的开发者建议从 csrc 中的现有算子入手遵循「实现 → 声明 → 注册 → 构建 → Python 导出 → 测试基准」六步流程并在提交前使用make format统一代码风格、运行make test与基准脚本验证正确性和性能。【免费下载链接】sglangSGLang is a high-performance serving framework for large language models and multimodal models.项目地址: https://gitcode.com/GitHub_Trending/sg/sglang创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考

sglang-kernel（sgl-kernel）内核库完全指南：安装构建、新增算子开发流程与体积分析

相关推荐

Awesome Copilot 实战：用 APPSYNC_JS 运行时构建生产级 AWS AppSync Event API 处理器（onPublish/onSubscribe 全指南）

Carbon 2021 路线图解读：从实验到可执行证据的加速路径

草莓腐烂二分类模型：轻量CNN实战与产线部署

更多资讯

STM32F407+μC/OS-II实现EtherCAT主站（DP83848单端口模拟双端口）

商超GEO项目实战：从数据采集到经营闭环

模型预测控制MPC入门：原理、实现与工程选型指南

CANN/GE图引擎获取输出索引API

5年测试经验被裁后自救：从功能测试到接口、自动化与性能

Thunderbird开源邮件客户端：功能解析与高级应用

今日更新

CopilotKit Vue Demo 与 React Demo 的 Parity 实现指南：路由、运行时端点与交互语义的全框架对齐

Harbor RBAC 实战：db_auth 模式下非管理员用户管理项目成员的完整验证与权限模型解析

STM32平衡车纯平衡程序从拆包到参数整定的完整调试指南

本周精选

超人会飞不算本事：系统稳定依赖清晰规则与边界设计

超人VS蜘蛛侠：拆解超级IP的影响力与传播方法论

基于CNN的调制信号识别：MATLAB实现时频图分类实战

本月热门

基于YOLO的焊缝缺陷检测：从数据标注到工业部署全流程实践

分布式服务的决策记录

创业产品的首版范围

读完文章，想聊聊您的网站？