常见问题(FAQ)
本 FAQ 侧重具体的技术问题和实操疑点,仓库定位、样例列表等基础信息请参见 README.md。
算子开发
我开发了一个新算子,目录结构和代码规范上需要注意什么?
参考 template/ 下的四个模板文件,分别对应算法说明、参考文献、算子代码框架、测试架构。特别注意:
- 算子目录名统一使用小写+下划线,如
my_operator/ - Host 侧 API 命名遵循
aclnn{算子名}约定 - Kernel 入口函数使用
__aicore__修饰符 - Tiling 参数需显式定义结构体,不能硬编码
- 所有数学函数(
sin、cos、sqrt等)需使用位运算 + Newton-Raphson 迭代替代,不能依赖<math.h>
Ascend C 算子中能否调用标准数学库(如 <math.h>)?
不能。Ascend C Kernel 侧不直接支持 <math.h> 和 __builtin_* 函数。所有数学函数需自行实现,推荐使用位操作 + Newton-Raphson 迭代(2 次迭代即可达到约 1×10⁻⁷ 精度)。已有 GAFF2、SHAKE 等算子的实现可作为参考。
PyTorch 算子与 Ascend C 算子在此仓库中如何共存?
两种形态均可提交。当前仓库以 Ascend C 自定义算子为主(位于 op_host/、op_kernel/),但每个算子配套的 Python 测试(tests/test_*.py)均基于 PyTorch 实现了参考计算逻辑,可直接作为 PyTorch → CANN 迁移的参考。如果你希望贡献纯 PyTorch 的模型样例(不涉及 Ascend C Kernel),放在对应子目录下并附上 README 即可。
我需要为算子提供哪些测试?
至少提供两类测试:
- C++ 单元测试(
tests/ut/op_kernel/test_*.cpp):与 CPU 参考实现逐结果对比 - Python 集成测试(
tests/test_*.py):端到端调用,对比 NumPy 和 PyTorch 结果
性能基准测试(tests/benchmark_*.py)虽然不是必须,但非常推荐。
架构差异
Ascend NPU 与 NVIDIA GPU 的计算架构有何本质区别?
NVIDIA GPU 采用 SIMT(Single Instruction, Multiple Thread)架构,而 Ascend NPU 采用 SIMD(Single Instruction, Multiple Data)为主的架构。这是两种截然不同的并行计算模型:
| 维度 | NVIDIA GPU (SIMT) | Ascend NPU (SIMD) |
|---|---|---|
| 执行模型 | 数千个轻量线程并行,每个线程有独立 PC | 向量化指令驱动,数据级并行(Vector/Cube Unit) |
| 线程分歧 | 允许(warp 内 divergence 会导致序列化) | 不存在线程概念,天然无分歧 |
| 编程思维 | 思考 "N 个线程各处理 1 个元素" | 思考 "一条指令处理 N 个元素" |
| 数据搬运 | 隐含在 Load/Store 指令中 | 显式搬运(GM → Local Memory → 计算 → 写回),通过 Pipe 机制流水 |
| 适合场景 | 不规则并行(稀疏、图、动态控制流) | 规则并行(矩阵乘法、卷积、向量化计算) |
| 开发语言 | CUDA | Ascend C |
这对算子开发有什么实际影响?
对于从 CUDA 迁移到 Ascend C 的开发者,需要转变几个核心思维:
-
不要逐元素编程:CUDA 中你写
kernel<<<grid, block>>>(data)启动 N 个线程各算一个元素;Ascend C 中你操作的是向量寄存器(如LocalTensor<float>),一条指令处理 64 或 128 个元素。 -
数据搬运需显式管理:Ascend C 的
DataCopy/Pipe机制要求开发者手动控制 GM(全局内存)→ Local Memory(片上缓存)的数据流动,这与 CUDA 中相对透明的 cache 层级不同。 -
Tiling 是关键:由于 Local Memory(约 2MB)远小于显存,大数组需要切分成 tile 逐块处理。本仓库的 LJ、GAFF2、DPD 等算子的
tiling文件是很好的参考。 -
向量化为王:尽量将标量运算改写为向量运算。例如计算
y = 1.0 / sqrt(x),不要写循环逐元素处理,而应调用Reciprocal+Sqrt向量指令。
Ascend C 的 Pipe 机制是什么?
Pipe 是 Ascend C 提供的数据搬运-计算流水线机制,允许数据加载、计算、存储三个阶段重叠执行:
┌─────────┐ ┌─────────┐ ┌─────────┐
│ 数据加载 │ ──→ │ 计算 │ ──→ │ 结果回写 │
│ (GM→LM) │ │ (Vector │ │ (LM→GM) │
│ │ │ /Cube) │ │ │
└─────────┘ └─────────┘ └─────────┘
↑ ↑ ↑
└───────── 流水线并行 ──────────────┘
通过 TPipe 和 TQue(队列)管理不同阶段的 buffer,可实现接近理论峰值的计算吞吐。当前 AI4MD 和 AI4PDE 的多个算子均利用了 Pipe 机制进行性能优化。
计算精度
NPU 上 float32 精度足够做分子动力学模拟吗?
已有算子的实测数据显示,在 Ascend C 上采用 float32 配合位运算数学函数实现,相对误差约 1×10⁻⁷,对 MD 模拟可忽略。SHAKE 约束在 float32 下经过 10 步 VV+SHK 循环后,坐标最大偏差约 9.54×10⁻⁶ nm(相对误差 < 0.01%)。
但对于需要长期能量守恒的模拟(如 NVE 系综),建议对照 CPU double 结果做更长时间的稳定性验证。
部分科学计算需要 FP64 双精度,但 Ascend 910B/910C 只支持到 FP32,怎么办?
这是一个关键限制,需要开发者充分了解:
硬件现状:
- Ascend 910B / 910C 的 AI Core(Cube Unit + Vector Unit)原生仅支持 FP32,没有像 NVIDIA GPU(V100 起支持 FP64 Tensor Core,A100 双精度达到 9.7 TFLOPS)那样的硬件 FP64 算力
- Ascend 950 系列在架构设计上是否有 FP64 支持,需以华为官方发布的规格为准
- 部分科学计算场景(如 DFT 中的对角化、MD 中的长时间积分、PDE 中的刚性方程求解)在 FP32 下可能出现精度不足
影响评估:
| 场景 | FP32 风险 | 建议 |
|---|---|---|
| 短时 MD 能量/力预测(< 1 ns) | 低(实测误差 < 1×10⁻⁶) | 可直接使用 |
| 长期 NVE 系综(> 10 ns) | 中(能量漂移可能累积) | 需与 FP64 基准对比验证 |
| 结构优化 / 过渡态搜索 | 中高(梯度精度影响收敛) | 建议用 CPU FP64 做最终验证 |
| 高精度电子结构计算 | 高(本征值精度不够) | 建议使用 CPU/GPU FP64 方案 |
| AI for PDE 逆问题 | 低(网络本身精度受限) | 可直接使用 |
实用建议:
- 场景决定精度需求:如果你的场景只需要"相对正确"的趋势(如 AI 势函数训练、材料筛选排序、工艺参数优化),FP32 完全够用
- 必须做精度验证:在切换到 NPU 之前,用 CPU double 结果作为基准,在你关注的全部物理量上做逐点对比
- 混合精度策略:可在 Host 侧用 CPU double 做累积误差敏感的部分(如长时间积分),NPU 上做大规模并行计算部分
- 社区交流:如果你在 FP32 下遇到了精度问题,欢迎在讨论区分享你的场景,也欢迎有 FP64 需求的用户一起推动 CANN 工具链的改进
力场算子的力约定是什么?
所有力的定义为 F=−∇EF = -\nabla E(指向能量降低方向),与 GROMACS 约定一致。返回力单位为 kJ/(mol⋅nm)\text{kJ}/(\text{mol} \cdot \text{nm})。
部署与集成
这些算子能否直接替换 GROMACS/LAMMPS 中的对应计算模块?
当前算子以 Ascend C 自定义算子的形式提供,提供了 aclnn 风格的 C++ API,可以通过 ACL(Ascend Computing Language)集成到自定义 MD 框架中。但尚未提供直接的 GROMACS/LAMMPS patch,如需在标准 MD 包中使用,需自行编写中间适配层。
算子如何在多卡或多节点上使用?
当前算子的实现是单设备(单 NPU)粒度的。如果需要多卡或多节点并行,需要在外部框架层(如自定义的 MPI + Python 调度)实现数据拆分和通信。CANN 提供了 HCCL(集合通信库)可用于多设备场景,但本仓库暂未提供多卡版的算子样例。
算子能否与 PyTorch 的 autograd 联动?
可以。你可以在 PyTorch 中通过 torch.ops 或自定义 torch.autograd.Function 封装 Ascend C 算子的 ACL 调用,从而在前向推理中利用 NPU 加速,同时保留 PyTorch 的自动微分能力。这对于需要可微分 MD、PINN 等场景非常有用。
性能分析
在 Atlas A2/A3 上运行这些算子,性能能达到什么水平?
以 Lennard-Jones 力场算子为例(Atlas A3 环境):
| 原子数 | CPU (PyTorch) | NPU 融合算子 | 加速比 |
|---|---|---|---|
| 64 | 0.54 ms | 0.57 ms | 0.96× |
| 128 | 25.57 ms | 0.75 ms | 34.21× |
| 256 | 174.96 ms | 0.85 ms | 206.23× |
| 512 | 183.00 ms | 1.45 ms | 126.36× |
在工业规模的体系(数千至数万原子)上,融合算子将古典力场计算从 CPU 上的"瓶颈步骤"变为 NPU 上的"几乎免费步骤"。但注意,跨场景的性能数据差异较大,建议针对自己的体系做基准测试。
我应该如何对我的算子做性能分析?
每个算子的 tests/ 目录下提供了 benchmark_*.py 脚本,作为性能基准测试的起点。建议至少测试 3-5 个不同规模(如小/中/大/超大),并记录:
- NPU Kernel 执行时间(通过 ACL 的 Event 计时)
- Host-Device 数据拷贝时间
- 与 CPU 参考实现的加速比
- 与 PyTorch GPU(如有)的加速比
兼容性与限制
算子是否支持动态 shape?
当前已发布算子(LJ、GAFF2、PME、SHAKE、VV)均支持动态原子数输入(通过参数传递 numAtoms),但不支持动态的 batch 维度。如果需要批次处理,可在 Host 侧循环调度多个独立的 Kernel 任务。
Ascend 910B、910C 和 950 系列对本仓库算子的支持情况如何?
| 硬件系列 | 产品型号 | 与本仓库算子的兼容性 | 说明 |
|---|---|---|---|
| Atlas A2 | Ascend 910B | ✅ 完全支持 | 当前开发和测试的主力平台,SOC_VERSION 设为 Ascend910B1 即可 |
| Atlas A3 | Ascend 910C | ✅ 完全支持(已验证) | LJ 力场算子的性能数据即在 Atlas A3 上测得,CMake 中指定对应 SOC_VERSION |
| Ascend 950 系列 | 待发布 | ❓ 待验证 | 作为昇腾下一代架构,预计提供更高算力和新指令集。算子级兼容性需等待硬件就绪后测试 |
| Atlas 200/500 A2 | Ascend 310B | ⚠️ 部分支持 | 推理卡,算力受限。小规模体系可运行,大规模 MD 模拟建议使用 A2/A3 |
| Atlas 800T A2 | Ascend 910B (训练) | ✅ 完全支持 | 训练服务器场景,与 A2 训练卡算子兼容 |
关键注意事项:
- SOC_VERSION 必须匹配:编译时需通过
-DSOC_VERSION=指定正确的芯片型号,不同系列使用的指令集和编译器优化策略不同 - 性能预期差异:同一算子在 910B 和 910C 上的性能可能不同(缓存大小、频率差异),建议在目标硬件上重新做基准测试
- 950 系列前瞻:虽然当前算子尚未在 950 上验证,但 Ascend C 作为昇腾统一的算子开发语言,具备良好的前向兼容性。后续硬件就绪后,本仓库会及时跟进适配
- 驱动和 CANN 版本要求:
- 910B:CANN ≥ 7.0,驱动 ≥ 23.0
- 910C:CANN ≥ 7.0.RC1,驱动 ≥ 24.0
- 具体版本兼容性矩阵请查阅华为官方文档
是否支持混合精度计算?
当前算子全部采用 float32 精度计算。CANN 工具链支持 FP16/INT8 量化,但科学计算场景对精度要求较高,暂未启用混合精度。未来规划中会评估在特定环节(如非键力计算中的截断判断、PME 倒易空间 FFT)引入半精度加速的可能性。
算子是否支持 Windows 开发环境?
算子本身在昇腾 NPU 上运行,开发环境需要在 Linux 服务器上配置 CANN 工具链。Windows 可用于代码编辑和文档编写,但编译和运行测试需要在 Linux 环境下完成。建议通过远程开发(VS Code Remote / SSH)方式连接到 Atlas 服务器。
贡献与交流
我发现的 bug 或想提出的功能建议应该发在哪里?
我想贡献的算子不在 roadmap 上,可以提交吗?
可以。roadmap 列出的是团队优先关注的方向,但社区贡献不受 roadmap 限制。只要算子符合"材料/化工/钢铁/油气等工业场景中的模拟仿真或预测优化"这一主题,且提供了完整文档和测试,即可提交 PR。建议先在讨论区发起话题,获取反馈后再开始编码。
我如何在自己的文章中引用这个仓库或其中的算子?
建议引用格式:
mat-chem-sim-pred: CANN-based operator library for material and chemical simulation and prediction.
https://gitcode.com/cann/mat-chem-sim-pred
引用特定算子时,附上该算子的算法说明文档链接(docs/algorithm.md)和对应的 DOI(如有)。
维护团队如何分工?
- Maintainer:负责仓库整体方向、SIG Review 和合入决策
- Committer:负责日常代码审查、Issue 回复和 CI 维护
如有紧急问题,可直接 @Maintainer 列表中的成员。