质量与性能的基石：CANN算子开发中的调试、测试与验证深度实践

摘要：

在高性能 AI 计算领域，一个算子（Operator）的价值不仅仅在于其实现了特定功能，更在于其正确性、鲁棒性与高效性。在华为 CANN（Compute Architecture for Neural Networks）架构中，由于计算发生在 NPU（神经网络处理单元）上，传统的 CPU 调试手段往往鞭长莫及，这使得算子的调试与验证成为开发流程中的关键挑战。

昇腾训练营报名链接： https://www.hiascend.com/developer/activities/cann20252#cann-camp-2502-intro

训练营简介： 2025 年昇腾 CANN 训练营第二季，基于 CANN 开源开放全场景，推出 0 基础入门系列、码力全开特辑、开发者案例等专题课程，助力不同阶段开发者快速提升算子开发技能。获得 Ascend C 算子中级认证，即可领取精美证书，完成社区任务更有机会赢取华为手机，平板、开发板等大奖

声明： 该文仅是为了记录 CANN 训练营的学习过程所用，不参与任何商业用途

一、 调试的艺术：TKC++孪生调试技术

在 NPU 上直接调试（Debug）核函数（Kernel）是极其困难的。CANN 框架为此提供了一套巧妙的孪生调试机制：同一套算子代码可以在 CPU 模式下进行模拟运行和精度调试，然后再无缝切换到 NPU 模式下执行。这极大地降低了调试门槛

1. GDB：CPU模式下的断点调试

当算子在 CPU 模式下运行时，我们可以像调试普通 C++ 程序一样，使用 GDB 进行单步跟踪和变量检查。

操作流程示例：

1. 编译CPU模式：在编译脚本中指定CPU模式，生成可执行文件（例如add_tik2_cpu）。
2. 启动GDB：gdb ./add_tik2_cpu
3. 设置断点：(gdb) b add_tik2_kernel（假设核函数名为add_tik2_kernel）
4. 运行：(gdb) r
5. 调试：程序将在核函数入口处暂停，此时可以使用 n (next), s (step), p (print) 等标准 GDB 命令检查代码逻辑和变量状态。

这种方式对于排查复杂的算法逻辑错误（如偏移量计算、循环边界等）至关重要。

2. 打印调试：NPU模式的“唯一窗口”

虽然 GDB 功能强大，但它仅限于 CPU 模式。在 NPU 模式下，printf 或 std::cout 语句是无效的。为了在开发阶段获取必要信息，我们必须使用 CANN 提供的特定宏 CCE_KT_TEST 来隔离打印语句，使其仅在 CPU 模式下编译和执行

代码示例：在核函数中安全地添加打印

#include "kernel_utils.h" // 确保包含了CANN工具头文件

extern "C" __global__ __aicore__ void add_tik2_kernel(
    __gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z,
    __gm__ uint8_t* tilingData) 
{
    // ... 解析 Tiling ...
    GET_TILING_DATA(tiling, tilingData, AddTik2TilingData);

    // 仅在CPU孪生调试模式下编译和执行
    #ifdef CCE_KT_TEST
    // 打印当前核心(Core)ID和它需要处理的数据块数
    printf("Debug Info: CoreID = %u, TileNum = %u\n", GetBlockIdx(), tiling.tileNum);
    #endif

    // ... 核心业务逻辑 ...
}

最佳实践：始终使用 CCE_KT_TEST 宏来包裹调试用的打印语句。这能确保您的代码在CPU模式下（如执行UT时）可以输出调试信息，同时在编译NPU模式（用于ST或部署）时，这些语句会被自动剔除，不会引入性能开销或编译错误。

二、 质量的守门人：UT与ST测试框架

CANN 算子开发推崇“测试驱动”的理念。一套完备的算子交付物必须包含相应的单元测试（UT）和系统测试（ST）

对比维度	单元测试 (Unit Testing, UT)	系统测试 (System Testing, ST)
测试目标	验证**核函数（Kernel）的算法逻辑	验证算子全链路**（Host+Device）的功能
运行环境	CPU模拟环境，不依赖NPU硬件	真实NPU硬件（如Ascend 910/310）
关注点	算法的逻辑正确性	端到端功能、NPU执行精度、Host侧交互
核心工具	GTest框架（C++）	msopst 工具（Python+JSON）
执行命令	build.sh -u tik2	bash run_case.sh

三、 单元测试 (UT) - 聚焦核函数逻辑

UT 的目的是在不依赖 NPU 硬件的情况下，快速验证核函数本身的算法是否正确

关键点：UT的Tiling解析

在 UT 环境中，tilingData 同样需要被解析。但与真实 NPU 环境不同，UT 中的 tiling.h 定义只保留了 CPU 模式下的处理逻辑，即将传入的Tiling指针直接转换为结构体指针，以便测试代码可以方便地“伪造” Tiling 数据

UT测试步骤：

1. 添加实现文件：在 cann_op_contrib/community/ops/ 目录下添加算子实现。
2. 添加测试文件：在 cann_op_contrib/community/tests/ut/ 目录下新建用例目录，并编写GTest测试代码。
   • 在GTest代码中，开发者会手动创建输入数据、手动填充Tiling结构体，然后调用核函数（在CPU上模拟执行），最后 ASSERT_EQ 检查输出结果是否符合预期。
3. 编译并执行：在 cann_op_contrib 目录下执行 ./build.sh -u tik2。

UT 是保证核函数算法鲁棒性的第一道防线，它运行速度快，适合在开发过程中频繁执行

四、 系统测试 (ST) - 真实硬件的全链路验证

ST 是算子交付前的最后一道关卡。它不再是模拟，而是在真实NPU硬件上运行算子，并将其输出与“黄金标准”（通常是Numpy的计算结果）进行对比，以验证其端到端的功能和精度。

CANN 为此提供了强大的 msopst 测试工具，其流程由三个核心文件驱动：

1. 核心文件一：add_tik2.json (测试用例定义)

此 JSON 文件定义了 ST 要执行的测试场景，包括算子名称、输入输出的 Shape、数据类型（DType）等

{
  "op_name": "AddTik2",
  "input_desc": [
    { "name": "x", "shape": [8, 1024], "type": "float32" },
    { "name": "y", "shape": [8, 1024], "type": "float32" }
  ],
  "output_desc": [
    { "name": "z", "shape": [8, 1024], "type": "float32" }
  ],
  "precision_standard": "rtol=1e-3,atol=1e-3"
}

2. 核心文件二：test_add_tik2_data.py (黄金数据生成)

msopst 工具会根据JSON文件生成随机的输入数据，然后调用此Python脚本来计算出理论上的正确答案（黄金数据）。

import numpy as np

def get_golden_data(input_data_list):
    """
    使用Numpy计算Add算子的期望输出
    input_data_list[0] 对应 "x"
    input_data_list[1] 对应 "y"
    """
    input_x = input_data_list[0]
    input_y = input_data_list[1]
    
    # 使用Numpy的标准实现作为黄金标准
    golden_data = np.add(input_x, input_y)
    
    # ST框架要求返回一个list
    return [golden_data]

3. 核心文件三：run_case.sh (执行脚本)

这是一个封装了 msopst 命令的Shell脚本，它会自动完成以下所有工作：

执行bash run_case.sh后，如果st_report.json中显示"result": "Success"，则代表算子通过了真实硬件的考验。

五、 性能调优：超越“正确”，追求“高效”

功能和精度验证通过后，我们还需要回答一个问题：这个算子跑得快吗？

CANN 提供了性能采集工具 msprof，它可以收集算子在 NPU 上执行时的详细硬件性能数据。

性能采集步骤：

1. 设置环境变量：source /usr/local/Ascend/ascend-toolkit/set_env.sh

2. 编译NPU可执行文件：确保算子被编译为NPU模式（例如 bash run.sh add_tik2 ascend910 Aicore npu）。

3. 使用msprof采集：启动 msprof 来分析NPU可执行文件。

   # 示例：采集AI Core的管线利用率(pipeUtilization)等指标
   msprof --application="./add_tik2_npu" \
          --output="./profiling_data" \
          --ai-core=on \
          --aic-metrics="pipeUtilization"
   

4. 分析报告：msprof 会在 profiling_data/ 目录下生成一系列CSV报告，例如 op_summary_0_1.csv。

通过分析这些指标，开发者可以判断算子是否存在瓶颈（例如 AI Core 利用率低、内存等待时间长等），并反过来指导Tiling策略的优化或核函数代码的重构

六、总结

CANN 算子的开发绝非“一锤子买卖”。本文从“质量保证”的视角出发，详细拆解了从开发初期的孪生调试（GDB与Printf），到中期的功能验证（UT与ST），再到后期的性能验证（msprof）的全链路流程

一个高质量的 AI 算子，是功能正确性、高精度和高性能的结合体。只有建立起这样一套严谨的调试与验证闭环，开发者才能构建出真正健壮、高效、可信赖的 AI 底层算子库，为上层 AI 模型的稳定运行提供坚实的基石