forked from ccf-ai-infra/Intro-ops
docs: add Chinese README and operator guide
Co-authored-by: wawahejun <hejunlbbc@gmail.com>
This commit is contained in:
yutianyu
parent
aec1073394
commit
1368fa2886
29
README.md
29
README.md
|
|
@ -3,9 +3,9 @@
|
|||
This repository is a training-oriented GPU operator runtime. It is intentionally
|
||||
small, but its workflow mirrors production operator libraries:
|
||||
|
||||
1. Define the operator contract in `ops/<op>/operator.yaml`.
|
||||
2. Generate build and runtime registries from the manifest.
|
||||
3. Implement a backend-specific descriptor lifecycle.
|
||||
1. Create a directory under `ops/<op>/` with backend implementations.
|
||||
2. The build system auto-discovers sources by directory convention.
|
||||
3. Implement a backend-specific descriptor lifecycle (create, workspace, execute, destroy).
|
||||
4. Expose a Python API with out-of-place, out-variant, and prepared execution.
|
||||
5. Validate correctness against PyTorch and benchmark steady-state execution.
|
||||
|
||||
|
|
@ -37,23 +37,34 @@ cmake --build . -j$(nproc)
|
|||
## Validate
|
||||
|
||||
```bash
|
||||
python tools/validate_operator_manifest.py --ops-root ops --tests-root tests
|
||||
python tests/run_ops.py --op copy --backend nvidia --mode all
|
||||
CAMP_BUILD_DIR=build pytest tests/ -v --backend nvidia
|
||||
pytest tests/ -v --backend tilelang
|
||||
python tests/bench_all.py --backend nvidia --profile tests/perf_profiles/local_gpu.yaml
|
||||
python tests/run_ops.py --op all --backend nvidia --mode bench
|
||||
```
|
||||
|
||||
The TileLang backend requires the `tilelang` Python package.
|
||||
The TileLang backend requires the `tilelang` Python package.
|
||||
|
||||
## Adding a New Operator
|
||||
|
||||
1. Create `ops/<name>/nvidia/<name>_cuda.h` with the C API (4 functions: create, workspace, execute, destroy).
|
||||
2. Create `ops/<name>/nvidia/<name>_cuda.cu` with the CUDA implementation.
|
||||
3. Create `python/operator_runtime/ops/<name>.py` using `ctypes_bindings.bind_*` functions.
|
||||
4. Create `tests/cases/<name>.py` with `correctness_cases()`, `api_error_cases()`, and `benchmark_cases()`.
|
||||
5. Create `tests/ops/test_<name>.py` and `tests/bench/<name>.py`.
|
||||
6. Re-run `cmake ..` in the build directory (the glob will pick up the new `.cu` file).
|
||||
7. Register the public API in `python/operator_runtime/__init__.py`.
|
||||
|
||||
No YAML, no code generation, no registration step.
|
||||
|
||||
## Production Mapping
|
||||
|
||||
| Training concept | Production equivalent |
|
||||
| --- | --- |
|
||||
| `operator.yaml` | reviewed operator spec / manifest |
|
||||
| directory convention `ops/<op>/nvidia/*.cu` | build system auto-discovery / operator registry |
|
||||
| C header `ops/<op>/nvidia/<op>_cuda.h` | reviewed operator API contract |
|
||||
| descriptor lifecycle | create, workspace, execute, destroy |
|
||||
| generated registry | operation table / backend registry |
|
||||
| `tests/cases/<op>.py` | correctness, layout, and API contract coverage |
|
||||
| `PerformanceResult` | profiler report row with latency, bytes, flops, bandwidth |
|
||||
| eager TileLang kernel | puzzle-stage kernel using `T.empty(...)` return values |
|
||||
| lazy TileLang `out_idx` template | TileOPs-style kernel factory and output-position contract |
|
||||
| lazy TileLang `out_idx` template | TileOPs-style kernel factory and output-position contract |
|
||||
|
|
|
|||
|
|
@ -0,0 +1,56 @@
|
|||
# Intro-ops训练营
|
||||
|
||||
这个仓库是一个面向训练的 GPU 算子运行时。它体积很小，但工作流尽量贴近真实的算子库开发流程：
|
||||
|
||||
1. 在 `ops/<op>/` 下创建后端实现目录。
|
||||
2. 构建系统通过目录约定自动发现源码。
|
||||
3. 实现后端专属的 descriptor 生命周期（create、workspace、execute、destroy）。
|
||||
4. 提供 Python API，支持 out-of-place、out-variant 和 prepared 执行。
|
||||
5. 通过 PyTorch 做正确性验证，并做稳定态性能测试。
|
||||
|
||||
## 算子
|
||||
|
||||
| 算子 | NVIDIA C++ | TileLang | MetaX |
|
||||
| --- | --- | --- | --- |
|
||||
| `copy` | 可运行 | 安装 TileLang 后可运行 | stub |
|
||||
| `vector_add` | 可运行 | 安装 TileLang 后可运行 | stub |
|
||||
| `reduce_sum` | 可运行，row-wise fp32 | 安装 TileLang 后可运行 | stub |
|
||||
| `softmax` | 可运行，row-wise fp32 | 安装 TileLang 后可运行 | stub |
|
||||
|
||||
## 安装
|
||||
|
||||
```bash
|
||||
pip install -r requirements.txt
|
||||
```
|
||||
|
||||
## 构建
|
||||
|
||||
```bash
|
||||
mkdir -p build
|
||||
cd build
|
||||
cmake .. -DCAMP_ENABLE_NVIDIA=ON -DCAMP_ENABLE_METAX=OFF
|
||||
cmake --build . -j$(nproc)
|
||||
```
|
||||
|
||||
## 验证
|
||||
|
||||
```bash
|
||||
python tests/run_ops.py --op copy --backend nvidia --mode all
|
||||
CAMP_BUILD_DIR=build pytest tests/ -v --backend nvidia
|
||||
pytest tests/ -v --backend tilelang
|
||||
python tests/run_ops.py --op all --backend nvidia --mode bench
|
||||
```
|
||||
|
||||
TileLang 后端需要安装 `tilelang` Python 包。
|
||||
|
||||
## 生产映射
|
||||
|
||||
| 训练概念 | 生产等价物 |
|
||||
| --- | --- |
|
||||
| `ops/<op>/nvidia/*.cu` 目录约定 | 构建系统自动发现 / 算子注册 |
|
||||
| `ops/<op>/nvidia/<op>_cuda.h` 头文件 | 经过评审的算子 API 契约 |
|
||||
| descriptor 生命周期 | create、workspace、execute、destroy |
|
||||
| `tests/cases/<op>.py` | 正确性、布局和 API 契约覆盖 |
|
||||
| `PerformanceResult` | 包含延迟、字节数、FLOPs、带宽的 profiler 报表行 |
|
||||
| eager TileLang kernel | 使用 `T.empty(...)` 返回值的 puzzle-stage kernel |
|
||||
| lazy TileLang `out_idx` template | TileOPs 风格的 kernel factory 和输出位置契约 |
|
||||
|
|
@ -0,0 +1,208 @@
|
|||
# 如何开发一个新算子
|
||||
|
||||
本文只描述当前仓库里新增算子的实际流程，不展开代码细节。
|
||||
|
||||
## 目标
|
||||
|
||||
在当前项目里，新增一个可运行算子的最小闭环包括四部分：
|
||||
|
||||
1. 在 `ops/<算子名>/` 下补后端实现。
|
||||
2. 在 `python/operator_runtime/ops/` 下补 Python API。
|
||||
3. 在 `tests/` 下补正确性测试和 benchmark 入口。
|
||||
4. 重新构建并验证。
|
||||
|
||||
## 当前仓库的真实约束
|
||||
|
||||
- 现在没有 `operator.yaml`。
|
||||
- 现在没有代码生成步骤。
|
||||
- 现在没有单独的 operator registry 文件。
|
||||
- NVIDIA 的 `.cu` 文件通过目录约定自动发现。
|
||||
- 新增或重命名 `.cu` 后，需要重新执行一次 `cmake ..`。
|
||||
|
||||
也就是说，新增算子主要依赖目录结构和命名约定，而不是额外的注册配置。
|
||||
|
||||
## Step 1：明确算子接口
|
||||
|
||||
开始前先确认：
|
||||
|
||||
1. 这个算子有几个输入、几个输出。
|
||||
2. 输出 shape 是否和输入一致。
|
||||
3. 是否要求输入输出 dtype 一致。
|
||||
4. 是否只支持 contiguous tensor。
|
||||
5. 是否需要额外参数，例如 `dim`、`scalar`。
|
||||
6. 是否需要 workspace。
|
||||
|
||||
这一步的目的，是确定后面 C API 和 Python 绑定该按哪种模式实现。
|
||||
|
||||
## Step 2：创建算子目录
|
||||
|
||||
先在 `ops/<算子名>/` 下建立对应目录。
|
||||
|
||||
当前建议至少补齐：
|
||||
|
||||
- `ops/<算子名>/nvidia/`
|
||||
- `python/operator_runtime/ops/<算子名>.py`
|
||||
- `tests/cases/<算子名>.py`
|
||||
- `tests/ops/test_<算子名>.py`
|
||||
- `tests/bench/<算子名>.py`
|
||||
|
||||
如果后续要支持 TileLang 或 MetaX，再分别补 `tilelang/` 或 `metax/`。
|
||||
但对当前流程来说，NVIDIA 版本是最小必需项。
|
||||
|
||||
## Step 3：实现 NVIDIA 后端
|
||||
|
||||
`ops/<算子名>/nvidia/` 这一层负责 C++/CUDA 实现。
|
||||
|
||||
通常需要这几部分：
|
||||
|
||||
1. kernel 文件
|
||||
2. C API 头文件
|
||||
3. C API 实现文件
|
||||
|
||||
这里最关键的是保持现有命名约定一致，因为 Python 侧会按固定符号名去找函数。
|
||||
|
||||
一个 NVIDIA 算子需要完整提供四个生命周期接口：
|
||||
|
||||
1. create
|
||||
2. workspace
|
||||
3. execute
|
||||
4. destroy
|
||||
|
||||
整体流程可以理解为：
|
||||
|
||||
1. `create`：检查输入是否合法，并保存运行所需信息。
|
||||
2. `workspace`：返回执行需要的临时内存大小。
|
||||
3. `execute`：按 dtype 和参数启动 kernel。
|
||||
4. `destroy`：释放 descriptor。
|
||||
|
||||
## Step 4：补 Python 绑定
|
||||
|
||||
Python 入口放在 `python/operator_runtime/ops/<算子名>.py`。
|
||||
|
||||
当前项目对外通常暴露三类接口：
|
||||
|
||||
1. `prepare_<算子名>`
|
||||
2. `<算子名>_`
|
||||
3. `<算子名>`
|
||||
|
||||
职责分工一般是：
|
||||
|
||||
1. `prepare_<算子名>`：完成参数检查、绑定底层函数、创建 descriptor 和 workspace。
|
||||
2. `<算子名>_`：接收调用方提供的输出 tensor，执行一次。
|
||||
3. `<算子名>`：自动分配输出 tensor，再调用 `<算子名>_`。
|
||||
|
||||
如果你的算子签名和现有 unary、binary、reduce-like 模式一致，就沿用现有 helper。
|
||||
如果签名比较特殊，例如带额外标量参数，就单独写一层绑定。
|
||||
|
||||
## Step 5：导出到公共 API
|
||||
|
||||
新增 Python 文件后，还要把公开接口补到：
|
||||
|
||||
1. `python/operator_runtime/ops/__init__.py`
|
||||
2. `python/operator_runtime/__init__.py`
|
||||
|
||||
至少需要把这三个名字暴露出去：
|
||||
|
||||
1. `<算子名>`
|
||||
2. `<算子名>_`
|
||||
3. `prepare_<算子名>`
|
||||
|
||||
这样测试和用户代码才能直接导入。
|
||||
|
||||
## Step 6：补测试用例数据
|
||||
|
||||
`tests/cases/<算子名>.py` 只负责组织测试数据。
|
||||
|
||||
当前仓库一般分成三类：
|
||||
|
||||
1. `correctness_cases()`
|
||||
2. `api_error_cases()`
|
||||
3. `benchmark_cases()`
|
||||
|
||||
建议在这里把不同 shape、dtype、容差、错误场景、性能场景分开整理，避免把 case 直接写死在测试函数里。
|
||||
|
||||
## Step 7：补正确性测试
|
||||
|
||||
`tests/ops/test_<算子名>.py` 主要负责三件事：
|
||||
|
||||
1. 正确性对比
|
||||
2. API contract 检查
|
||||
3. prepared 执行复用检查
|
||||
|
||||
正确性测试通常是拿 PyTorch 结果做对照。
|
||||
API contract 测试主要覆盖 shape 不匹配、dtype 不匹配、非 contiguous 等场景。
|
||||
如果底层走 descriptor lifecycle，建议补一个 prepared 多次执行的测试，确认 descriptor 可以复用。
|
||||
|
||||
## Step 8：补 benchmark
|
||||
|
||||
`tests/bench/<算子名>.py` 负责性能入口。
|
||||
|
||||
当前流程里，一般会：
|
||||
|
||||
1. 从 `benchmark_cases()` 取输入规模。
|
||||
2. 构造 CUDA tensor。
|
||||
3. 测量自定义算子耗时。
|
||||
4. 测量对应 PyTorch 实现耗时。
|
||||
5. 汇总成性能结果。
|
||||
|
||||
这一步的目标不是做复杂分析，而是保证新算子已经接入仓库现有 benchmark 流程。
|
||||
|
||||
## Step 9：重新配置和编译
|
||||
|
||||
因为 `ops/CMakeLists.txt` 是通过 glob 自动发现 `ops/*/nvidia/*.cu`，所以新增 `.cu` 之后要重新配置。
|
||||
|
||||
顺序是：
|
||||
|
||||
1. 进入 `build/` 目录。
|
||||
2. 重新执行 `cmake ..`。
|
||||
3. 再执行编译。
|
||||
|
||||
如果你只改了 Python，不新增 `.cu`，通常不需要重新配置。
|
||||
但只要新加了 NVIDIA 源文件，就必须重新跑一次 CMake。
|
||||
|
||||
## Step 10：验证
|
||||
|
||||
建议按下面顺序验证：
|
||||
|
||||
1. 先跑单算子的 pytest。
|
||||
2. 再跑更大范围的算子测试。
|
||||
3. 最后跑 benchmark 或统一入口。
|
||||
|
||||
优先确认：
|
||||
|
||||
1. 输出结果和 PyTorch 一致。
|
||||
2. 错误输入能稳定报错。
|
||||
3. descriptor/workspace 流程正常。
|
||||
4. 没有破坏已有算子。
|
||||
|
||||
## 推荐开发顺序
|
||||
|
||||
如果你要新增的是一个普通算子，建议按这个顺序推进：
|
||||
|
||||
1. 先参考最接近的现有算子选模板。
|
||||
2. 先打通 NVIDIA 后端最小可运行版本。
|
||||
3. 再补 Python API。
|
||||
4. 再补 cases、pytest、benchmark。
|
||||
5. 最后补其他后端或做性能优化。
|
||||
|
||||
## 现有模板怎么选
|
||||
|
||||
- `copy`：适合最简单的单输入单输出流程。
|
||||
- `vector_add`：适合标准 elementwise 双输入算子。
|
||||
- `reduce_sum`：适合带 reduce 维度的算子。
|
||||
- `softmax`：适合带更明确 shape 约束和归一化逻辑的算子。
|
||||
|
||||
如果新算子本质上是普通 elementwise，优先参考 `vector_add`。
|
||||
|
||||
## 最终检查清单
|
||||
|
||||
提交前至少确认以下内容都已完成：
|
||||
|
||||
1. `ops/<算子名>/nvidia/` 已补齐实现。
|
||||
2. `python/operator_runtime/ops/<算子名>.py` 已补齐。
|
||||
3. 两个 `__init__.py` 已导出新接口。
|
||||
4. `tests/cases/<算子名>.py` 已补数据。
|
||||
5. `tests/ops/test_<算子名>.py` 已补测试。
|
||||
6. `tests/bench/<算子名>.py` 已补 benchmark。
|
||||
7. 新增 `.cu` 后已经重新执行过 `cmake ..`。
|
||||
8. 至少完成一次单算子验证。
|
||||
Loading…
Reference in New Issue