From 1368fa2886f5ce07cb98a6593d51efa68d213d94 Mon Sep 17 00:00:00 2001
From: yutianyu <SqCBeFObcpwfC3SvwsFmWI+yutianyu@noreply.cnb.cool>
Date: Mon, 4 May 2026 19:16:21 +0800
Subject: [PATCH] docs: add Chinese README and operator guide

Co-authored-by: wawahejun <hejunlbbc@gmail.com>
---
 README.md                      |  29 +++--
 README.zh.md                   |  56 +++++++++
 docs/how-to-add-an-operator.md | 208 +++++++++++++++++++++++++++++++++
 3 files changed, 284 insertions(+), 9 deletions(-)
 create mode 100644 README.zh.md
 create mode 100644 docs/how-to-add-an-operator.md

diff --git a/README.md b/README.md
index 99ea244..5cdd727 100644
--- a/README.md
+++ b/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 |
\ No newline at end of file
+| lazy TileLang `out_idx` template | TileOPs-style kernel factory and output-position contract |
diff --git a/README.zh.md b/README.zh.md
new file mode 100644
index 0000000..a670d7f
--- /dev/null
+++ b/README.zh.md
@@ -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 和输出位置契约 |
diff --git a/docs/how-to-add-an-operator.md b/docs/how-to-add-an-operator.md
new file mode 100644
index 0000000..92730d5
--- /dev/null
+++ b/docs/how-to-add-an-operator.md
@@ -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. 至少完成一次单算子验证。