forked from ccf-ai-infra/Intro-ops
112 lines
4.4 KiB
Markdown
112 lines
4.4 KiB
Markdown
# 常见问题解答 (FAQ)
|
||
|
||
> 最后更新：2026-06-05
|
||
|
||
## 环境配置
|
||
|
||
### Q: 需要什么硬件才能运行 intro-ops？
|
||
|
||
**A:** 需要 NVIDIA GPU（支持 CUDA）或沐曦 MetaX GPU。Intel 集显、AMD GPU（无 ROCm 适配）、Apple Silicon 均无法运行。
|
||
|
||
### Q: 没有 NVIDIA GPU 怎么办？
|
||
|
||
**A:** 可以考虑租用云端 GPU（AutoDL、恒源云等平台），租一张 T4 或 GTX 1060 即可跑通全部算子。
|
||
|
||
### Q: CUDA Toolkit 需要什么版本？
|
||
|
||
**A:** 建议 CUDA 11.8 或以上。CMake 用 `"native"` 架构参数可自动适配你的 GPU。
|
||
|
||
### Q: 必须用 conda 管理 Python 环境吗？
|
||
|
||
**A:** 不强制，但推荐。主要依赖是 PyTorch >= 2.0、pytest >= 7.0、cmake >= 3.22。TileLang 后端还需 `pip install tilelang`。
|
||
|
||
### Q: Windows 上能跑吗？
|
||
|
||
**A:** 项目设计为 Linux 环境（构建脚本为 bash，产物为 `.so`）。Windows 上建议用 WSL2。
|
||
|
||
---
|
||
|
||
## 编译构建
|
||
|
||
### Q: 构建报 `FindCUDA failed` 怎么办？
|
||
|
||
**A:** 检查 CUDA Toolkit 是否安装，`nvcc --version` 是否可运行，`CMAKE_CUDA_COMPILER` 是否正确。
|
||
|
||
### Q: 每次改完 kernel 都要重新构建吗？
|
||
|
||
**A:** NVIDIA 后端需要重新构建（`bash scripts/build_nvidia.sh build`）。TileLang 后端是 JIT 编译，改完 Python 代码直接跑测试即可。
|
||
|
||
### Q: `CMake Error: Unknown CUDA architecture` 怎么处理？
|
||
|
||
**A:** 修改 `CMakePresets.json`，将 `CMAKE_CUDA_ARCHITECTURES` 改为 `"native"` 或你的 GPU 对应架构编号。
|
||
|
||
### Q: 构建成功但 `import` 时报找不到 `.so` 文件？
|
||
|
||
**A:** 检查环境变量 `CAMP_BUILD_DIR` 是否指向正确的构建目录（如 `build-nvidia`）。
|
||
|
||
---
|
||
|
||
## Kernel 编写
|
||
|
||
### Q: 四个算子应该按什么顺序学习？
|
||
|
||
**A:** copy → vector_add → reduce_sum → softmax，难度递增。copy 最简单（纯内存搬运），softmax 需要理解数值稳定性和 online 算法。
|
||
|
||
### Q: grid-stride loop 为什么能处理任意大小的 tensor？
|
||
|
||
**A:** 每个线程在循环中处理多个元素（步长为 `gridDim.x * blockDim.x`），而不是只处理一个。这样无论总元素数是多少，只要循环条件 `idx < N` 就能覆盖。
|
||
|
||
### Q: `__syncthreads()` 什么情况下必须用？
|
||
|
||
**A:** 当 block 内线程通过 shared memory 交换数据时，写入 shared memory 后必须 `__syncthreads()` 确保所有线程都完成写入，之后才能安全读取。典型场景：reduce_sum 的树形归约每一步之间。
|
||
|
||
### Q: softmax 为什么要减最大值？
|
||
|
||
**A:** `exp(88.7) ≈ 1.6e38`，接近 FP32 上限 `3.4e38`。更大的输入值会导致 `exp()` 溢出为 inf。减最大值等价于分子分母同除 `exp(max)`，数学结果不变但数值稳定。
|
||
|
||
### Q: TileLang 和 CUDA kernel 需要功能完全一致吗？
|
||
|
||
**A:** 是的，两者应通过相同的测试用例。但实现方式不同——CUDA 控制线程级别，TileLang 控制 tile 级别。
|
||
|
||
---
|
||
|
||
## 测试与调试
|
||
|
||
### Q: 测试报 `tensor not close` 但肉眼看不出来？
|
||
|
||
**A:** 检查 `rtol`/`atol` 设置。FP16 建议 `rtol=1e-3`，FP32 建议 `rtol=1e-5`。如果使用 TileLang 后端，注意它用 `exp2`/`log2` 而非 `exp`/`log`，会产生细微差异。
|
||
|
||
### Q: 如何单独跑一个算子的测试？
|
||
|
||
**A:**
|
||
```bash
|
||
PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia pytest tests/op_tests/test_copy.py -v --backend nvidia
|
||
```
|
||
|
||
### Q: 如何同时跑正确性和 benchmark？
|
||
|
||
**A:**
|
||
```bash
|
||
PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia python tests/run_ops.py --op copy --backend nvidia --mode all
|
||
```
|
||
|
||
### Q: 测试结果不稳定怎么办？
|
||
|
||
**A:** 少量浮点误差波动（尤其是 reduce_sum 这类归约算子）是正常的。如果同一输入每次运行结果差异超过 `rtol=1e-5`（FP32）或 `rtol=1e-3`（FP16），检查是否有未初始化的内存或越界访问。
|
||
|
||
---
|
||
|
||
## 贡献流程
|
||
|
||
### Q: 我可以贡献什么？
|
||
|
||
**A:** 从 good-first-issue 标签入手比较合适。常见新手任务：修复文档错别字、补充测试用例、翻译文档、添加代码注释。详见 [how-to-submit-first-pr.md](how-to-submit-first-pr.md)。
|
||
|
||
### Q: PR 需要什么条件才能被合并？
|
||
|
||
**A:** 需要通过 CI（所有测试通过）、至少一位 reviewer 审核、代码风格符合规范、包含必要的测试。
|
||
|
||
### Q: 如何添加一个新算子？
|
||
|
||
**A:** 参考 [how-to-add-an-operator.md](how-to-add-an-operator.md)，完整流程包括 kernel 实现、C API、Python 绑定、测试和 benchmark。
|