目录
1相关名词解释
1.1. cudart（CUDA Runtime）
1.2. cuBLAS
1.3. CUTLASS
1.4. cuDNN
1.5. Triton（多半指 OpenAI Triton）
1.6. Marlin（vLLM 等推理场景里常见）
1.7. AITER（你仓库里和 ROCm 相关）
1.8 疑问：cutlass和cublas有什么区别
2 大模型量化
2.1 NVIDIA Model Optimizer（ModelOpt）
2.2 华为昇腾 msModelSlim
2.3 llm-compressor
2.4 对比
2.5 疑问：我们量化用什么
3 deepgemm
3.1 gf
3.2 bh
3.3
4 lightop
5 lmslim
6 什么是marlin
6.1 开源的marlin
6.2 我们内部所谓的marlin
和「跟 Marlin 有没有关系」怎么折中
7 aiter
7.1 开源的aiter
7.1.1 它到底是什么
7.1.2「做了哪些优化」——官方怎么说
7.1.3 用什么写的：不是「只有 HIP」，而是多条后端
7.1.4 文档里列的能力范围（便于你对照「是不是全家桶」）
7.2 我们的aiter
以下名词可能有理解错误的地方，仅仅是初步理解和整理的个人笔记，仅供参考。
1相关名词解释
1.1. cudart（CUDA Runtime）
| 项目 | 说明 |
|---|---|
常见写法 | CUDA Runtime，API 前缀多为 |
缩写含义 | CUDA = Compute Unified Device Architecture；Runtime 即运行时库 |
作用 | 在应用进程里提供相对「好用」的接口：分配设备内存、 |
层次 | 在用户代码与 CUDA Driver API 之间；典型栈：应用 → cudart → CUDA 驱动 → GPU |
1.2. cuBLAS
| 项目 | 说明 |
|---|---|
全称 | cu BLAS = CUDA Basic Linear Algebra Subprograms（对标 CPU 上的 BLAS） |
作用 | GPU 上的矩阵乘、向量运算等通用线性代数例行程序，高度优化 |
层次 | 建立在 CUDA（含 cudart）之上；很多框架和专用库会调用 cuBLAS 做 GEMM 等 |
1.3. CUTLASS
| 项目 | 说明 |
|---|---|
全称 | CUT LASS = CUDA Templates for Linear Algebra Subroutines Software（常见简称 CUTLASS） |
作用 | 用 C++ 模板拼出高性能 GEMM/卷积等；便于自定义 tiling、精度组合、融合算子 |
和 cuBLAS 关系 | 二者都在做「线性代数高性能」：CUTLASS 更偏可组合、可裁剪的源码级框架；新版 cuBLAS 内部也可能采用 CUTLASS 等技术路线实现部分路径 |
层次 | 依赖 CUDA；常被当作「写/生成 GEMM 家族内核的底层积木」 |
1.4. cuDNN
| 项目 | 说明 |
|---|---|
全称 | cu DNN = CUDA Deep Neural Network library |
作用 | 面向深度学习的原语：卷积、归一化、注意力相关、RNN 等，带多种算法选择与融合 |
和 cuBLAS 关系 | cuBLAS 偏通用 BLAS；cuDNN 偏 DL 专用算子与调度。内部实现常会用到 GEMM 等，可能与 cuBLAS/CUTLASS 路线协同 |
层次 | 同样在 CUDA/cudart 之上，通常位于「框架算子」之下的一层专用库 |
1.5. Triton（多半指 OpenAI Triton）
| 项目 | 说明 |
|---|---|
缩写 | 一般不作为短语缩写：它就是叫 Triton（类似编程语言名），不是「四个字母凑成的全称」 |
作用 | Python 风格写 GPU kernel → Triton 编译器生成 LLVM/NVPTX（CUDA）或 AMDGPU（ROCm）等后端代码 |
层次 | 与「手写 CUDA C++/CUTLASS」并列的一种内核开发方式：框架 / PyTorch 扩展 → Triton 源码 → 编译产物 → GPU。和 cuBLAS「直接调库」不是一类事情——Triton 是你自己写算法，编译器帮你 codegen |
1.6. Marlin（vLLM 等推理场景里常见）
| 项目 | 说明 |
|---|---|
缩写 | 常见用法里不作为标准技术缩写展开；项目名 Marlin（社区里的权重量化 GEMM 实现） |
作用 | 一类针对 权重量化（如 W4A16） 的高性能 GEMM/推理内核（你在 vLLM 里看到的 Marlin 路径即此类） |
层次 | 应用或推理引擎（如 vLLM）可选的一种专用内核实现，通常基于 CUDA；可能间接用到模板/汇编级优化，但不是像 cuBLAS 那样的通用数学库标准接口 |
1.7. AITER（你仓库里和 ROCm 相关）
| 项目 | 说明 |
|---|---|
名称 | AMD 在 ROCm 生态里的 AITER 库（你这边 |
缩写 | 官方材料里常见两种说法：「AI Tensor Engine for ROCm」，文档里也有 「AMD Inference and Training Enhanced Repository」 这类全称表述——指向同一套面向 ROCm 的算子/内核集合 |
作用 | 在 AMD GPU / ROCm 上提供 MoE、采样、注意力等优化算子（与 NVIDIA 栈里的 cuDNN/cuBLAS 部分功能「生态位」类似，但厂商与 API 不同） |
层次 | ROCm / HIP 运行时之上的 AMD 侧高性能 AI 算子库；和 cudart + cuBLAS/cuDNN 是不同硬件栈上的对应物，不是同一进程里并列的一层 |
1.8 疑问：cutlass和cublas有什么区别
我觉得cublas可以叫基础线性代数子程序库，cutlass线性代数子程序模板库，他们两个的区别就是，
如果用cublas，我们直接调用api去做线性代数或者说矩阵计算就好了，自己能左右的优化有限，或者说自己施展的空间有限，
如果用cutlass，那么我们可以自己去做一些优化，他更适合 深度定制（例如特殊的 epilogue、block 大小、融合）。
2 大模型量化
2.1 NVIDIA Model Optimizer（ModelOpt）
- 仓库：NVIDIA/Model-Optimizer
- 定位：NVIDIA 官方的统一模型优化库：量化（FP8/INT8/INT4/NVFP4 等）、剪枝、蒸馏、投机解码相关能力等，面向把模型压小、便于接到 TensorRT-LLM、TensorRT、vLLM、SGLang 等推理栈。
- 典型用法：在 PyTorch/Hugging Face 模型上做 PTQ 或配方化量化，再导出到下游框架。
- 心智：NVIDIA GPU 推理链路里的「离线优化一站式」，和 TensorRT-LLM / vLLM 强相关。
2.2 华为昇腾 msModelSlim
- 常见入口：GitCode 镜像如 Ascend/msmodelslim，文档如 msmodelslim.readthedocs.io
- 定位：昇腾（Ascend NPU）亲和的大模型压缩与量化工具：W8A8、W4A8、GPTQ、KV 量化、敏感层分析等（具体以版本文档为准）。
- 典型用法：把 LLM/MoE 等量化到昇腾可用格式，再接 MindIE、vLLM-Ascend 等。
- 心智：昇腾部署前的「自家 ModelOpt / 量化流水线」，硬件绑定 CANN / 昇腾，不是给 CUDA 用的通用包。
2.3 llm-compressor
- 背景：社区里常与 vLLM / Neural Magic 路线一起出现（仓库一般为 vllm-project/llm-compressor 或 Neural Magic 相关维护，以当前 GitHub 为准）。
- https://github.com/vllm-project/llm-compressor
- 定位：面向 LLM 的压缩与量化工具：GPTQ、AWQ、稀疏等，强调和 PyTorch / Hugging Face / vLLM 工作流贴近。
- 典型用法：在 GPU（多为 NVIDIA） 上对开源模型做量化，导出给 vLLM 等框架加载。
- 心智：开源社区向的「专用 LLM 量化工具」，生态上和 vLLM 绑定较深；不一定覆盖昇腾专用格式。
2.4 对比
| 项目 | 出品 / 生态 | 主要硬件联想 | 和推理框架关系 |
|---|---|---|---|
ModelOpt | NVIDIA | NVIDIA GPU | TensorRT-LLM、TensorRT、vLLM、SGLang |
msModelSlim | 华为昇腾 | 昇腾 NPU | MindIE、vLLM-Ascend 等 |
llm-compressor | 开源社区（vLLM 系） | 多为 NVIDIA GPU | vLLM、HF 模型 |
2.5 疑问：我们量化用什么
我们量化用的是llm-compressor。
我们的lmslim不是做量化的，而是一个量化后的算子库，相当于我们有些量化后的算子需要用lmslim，lmslim可以理解成跟deepgem， lightop，并列的东西。
3 deepgemm
3.1 deepseek的deepgemm
3.2 我们自己的deepgemm
应该也是基于开源的
3.2.1 gf
deepgemm是汇编实现的
deepgemm 是一个库，里面有很多接口， gemm 的接口是调用的 blaslt 库， 其他的咱们自己写的，对于gemm算子来说，hipblaslt库---->汇编---->cu文件---->py文件，
实现是和marlin强绑定的， marlin其实就是重排了下权重，更利于访存。
3.2.2 bh
hip和rocmblaslt都有，但主要是hip实现。
3.2.3
目前看，我们只用了deepgemm中的moe那块的gemm算子。
4 lightop
我觉得这个就是一个融合算子库，里面包含了几乎所有的算子。
这个主要就是用hip和汇编写的。
5 lmslim
我发现怎么这个lmslim里面又调用了lightop的东西，比如lmslim里面的fused_experts_impl_int8_marlin函数，调用了下面三个
moe_gemm_marlin_w8a8
fuse_silu_mul_quant
moe_gemm_marlin_w8a8
而这三个的定义都在lightop里面，
怎么感觉这个lmslim像是对算子的 一个封装，vllm可以调用lmslim然后进而调用到lightop。
6 什么是marlin
6.1 开源的marlin
https://github.com/IST-DASLab/marlin
6.2 我们内部所谓的marlin
那么这个函数的名字你看带marlin，但我感觉他的算子跟marlin算法没关系，我觉得他就是权重根据marlin重排了，所以这里算子也对应的 用这种方式实现，就是跟重排函数相匹配
- 「Marlin」在这里的作用，主要是 约定：权重已经按 Marlin 那套方式重排/打包（
w1/w2的 shape、block、scale 等和 Marlin W8A8 MoE 路径一致）。 moe_gemm_marlin_w8a8等名字表示：这是 吃「Marlin 排布权重」的 MoE W8A8 GEMM 实现（在你们是 Lightop 里），不是在数学上再实现一遍「Marlin 论文里的全部技巧」这个抽象概念。- 所以本质上是：数据布局叫 Marlin layout → 选/实现了与之匹配的算子；算子 = 和重排方式配套的内核家族，而不是「整个函数 = Marlin 论文算法本体」。
和「跟 Marlin 有没有关系」怎么折中
- 有关系：权重量化后的排布、scale 约定、以及为此写的 GEMM 沿用了 Marlin 风格 的推理路径。
- 没「整段照抄」那种关系：未必是上游仓库里同一份
marlin_cuda_kernel.cu；在你们栈上是 HIP/Lightop 等另一套实现，语义/布局对齐，代码可以完全不同。
一句话：名字里的 marlin 更适合理解成 「Marlin layout 上的 MoE INT8 路径」；你的说法——权重按 Marlin 重排，算子是为这种排布服务的实现——这样记 没问题。
7 aiter
7.1 开源的aiter
https://github.com/ROCm/aiter
7.1.1 它到底是什么
AITER（AI Tensor Engine for ROCm，文档里也写 AMD Inference and Training Enhanced Repository）是 AMD 在 ROCm 上维护的 高性能 AI 算子库：面向 推理 + 训练，给框架（vLLM、SGLang 等）提供 可直接调用的生产级 GPU 内核，带 Python / C++ API。
定位不是「某一个 GEMM」，而是 一整包算子（Attention、MoE、GEMM、Norm、量化、通信等）。
7.1.2「做了哪些优化」——官方怎么说
README 里强调的是 相对其它实现的加速倍数（具体场景依赖型号与形状），例如文档表格里的量级：
- MLA decode、MHA prefill 等大倍数场景
- Block-scaled Fused MoE、Block-scaled GEMM 等 数倍级
- 端到端 DeepSeek-R1（SGLang）等案例
这些数字背后是 算法与实现层面的组合，典型包括（通用术语，不限于某一代代码）：
| 方向 | 含义 |
|---|---|
Kernel 融合 | 把多处访存、多次 kernel launch 合成更少次数（例如 Norm+量化、MoE 里路由与计算的结合等）。 |
针对 CDNA 的矩阵指令 | 利用 MFMA 等矩阵乘加速路径，配合合适的 tile、寄存器与内存层级使用。 |
形状/配置的预调优 | 「Pre-tuned configs」——对常见 LLM 形状、batch、MI300/MI350 等 预先选好 block/tile，减少运行时摸索。 |
稀疏 / 块缩放等新特性 | 官方 News 里提到的 blockwise sparse Sage Attention、fused gated RMSNorm + group quantization、blockwise 相关 GEMM/MoE 等，属于 算子能力 + 数值格式的扩展。 |
Serving 路径 | Paged Attention、KV 相关优化，降低长上下文下的带宽与计算浪费。 |
通信 | Triton + Iris 做 GPU 侧发起的集合通信（README 单独提了 reduce-scatter、all-gather 等）。 |
注意：不是单一技巧，而是 不同算子选不同后端 + 融合 + 调参 + 针对 MI300/MI350 等发布的 tuned config。
7.1.3 用什么写的：不是「只有 HIP」，而是多条后端
README 明确列出三条内核后端：
- Triton（需 amd-triton，按 ROCm 版本选 wheel）
- Composable Kernel (CK) — AMD 的 C++ 模板化内核库
- Hand-tuned ASM — 手写汇编级调优
另外还有仓库内自带的开发基础设施：
- Opus（
csrc/include/opus/）：单头文件 C++ 模板，用来写 HIP 内核——封装 向量化 load/store、布局抽象、MFMA 包装 等，README 强调 缩短编译时间（相对「笨重」的扩展编译）。
可选依赖：
- FlyDSL：FusedMoE 里可用于混合精度 MoE；未安装时回退到 CK（README 原文）。
- Iris：Triton 侧 GPU 发起通信。
所以回答「是直接 HIP 还是调底层库」：
- 自定义 HIP：有（例如 Opus 路径下的 HIP 内核）。
- 调底层库：CK 本身就是 AMD 官方高性能 GEMM/卷积等内核组件库；Triton 走的是 amd-triton → ROCm 后端，最终在 GPU 上仍是 ROCm 工具链生成的代码。
- 汇编：部分热点用手写 ASM 再封装。
一句话：AITER = 多后端混合（Triton + CK + ASM + 自研 HIP 模板），不是「全盘只用裸 HIP 手写」或「全盘只调一个库」这一种。
7.1.4 文档里列的能力范围（便于你对照「是不是全家桶」）
官方文档归纳的典型类别：
- Attention：MHA、MLA、Paged Attention 等
- GEMM：FP16/BF16/FP8/INT4、预调优、t fusion
- MoE：融合 MoE、多种路由、量化专家权重
- Norm：RMSNorm、LayerNorm、融合变体
- 其它：RoPE、量化、逐元算子、通信（Triton/Iris）
硬件侧 README 表格主打 CDNA / Instinct（gfx90a、gfx942、gfx950 等 MI300/MI350 系列）。
7.2 我们的aiter
Aiter is an operator inheritance library derived from https://github.com/ROCm/aiter,
