6.1 调试技术

注：别忘了增加或者停止服务啊，一小时18块多。

6.1.1 日志输出调试

printf调试

最简单的调试方法就是打日志，在关键位置输出变量值，看看程序执行到哪里了，数据对不对。

Ascend C提供了printf API，可以在kernel里打印信息：

// 概念性示例
#include "ascendc.h"

extern "C" __global__ __aicore__ void MyKernel(...) {
    printf("Kernel started\n");
    
    int32_t block_id = GetBlockId();
    printf("Processing block %d\n", block_id);
    
    // 打印一些关键变量
    printf("Input size: %d\n", input_size);
    printf("Output size: %d\n", output_size);
    
    // 打印数据值（注意不要打印太多，影响性能）
    Scalar<float> first_value;
    first_value.Load(local_input, 0);
    printf("First value: %f\n", first_value.GetValue());
}

使用技巧：打印日志要注意，不要打印太多，会影响性能。只在关键位置打印，或者用条件编译控制。

DumpTensor调试

如果想看整个Tensor的内容，可以用DumpTensor API：

// 概念性示例
#include "ascendc.h"

extern "C" __global__ __aicore__ void MyKernel(
    GlobalTensor<float> input,
    GlobalTensor<float> output
) {
    // Dump输入Tensor
    DumpTensor(input, "input_tensor");
    
    // 进行计算
    Compute(...);
    
    // Dump输出Tensor
    DumpTensor(output, "output_tensor");
}

DumpTensor会把Tensor的内容保存到文件，可以在Host端查看。这对调试很有用，但要注意文件大小，大Tensor会生成很大的文件。

条件打印

有时候只想在特定条件下打印，可以用条件判断：

// 概念性示例
if (debug_mode) {
    printf("Debug info: value = %f\n", value);
}

// 或者用宏控制
#ifdef DEBUG
    printf("Debug: %s\n", message);
#endif

6.1.2 断点调试

CPU侧调试

Ascend C支持CPU侧调试，可以在CPU上模拟NPU的行为，用GDB等调试器调试。

孪生调试：CPU侧调试也叫孪生调试，就是在CPU上模拟NPU的执行，方便调试。

# 概念性命令
# 在CPU模式下运行算子
cpu_run_kernel MyKernel input output

# 用GDB调试
gdb ./my_program
(gdb) break MyKernel
(gdb) run
(gdb) print variable_name

CPU侧调试的好处是速度快，不需要NPU硬件，可以用标准的调试工具。但只能验证逻辑，不能验证性能。

NPU侧调试

NPU侧调试是在实际的NPU硬件上调试，能看到真实的执行情况。

NPU侧调试比较复杂，通常用日志和DumpTensor，或者用专门的调试工具。

调试模式：有些工具支持NPU的调试模式，可以单步执行，查看寄存器值等。

6.1.3 性能分析工具

msprof工具

msprof是CANN提供的性能分析工具，可以分析算子的性能。

# 概念性命令
# 启动性能分析
msprof --application="your_app" --output=./prof_output

# 查看分析结果
msprof --export=on --output=./prof_output

msprof会收集算子的执行时间、内存使用、NPU利用率等信息，生成性能报告。

性能报告内容：性能报告通常包括算子执行时间、各阶段耗时、内存访问模式、NPU利用率等。

MetricsProf API

在代码里也可以用MetricsProf API标记要分析的代码段：

// 概念性示例
#include "ascendc.h"

extern "C" __global__ __aicore__ void MyKernel(...) {
    // 开始性能分析
    MetricsProfStart();
    
    // 要分析的代码
    Compute(...);
    
    // 结束性能分析
    MetricsProfStop();
}

这样msprof工具就能知道要分析哪段代码，生成更精确的报告。

TRACE打点

TRACE打点可以标记代码的执行阶段：

// 概念性示例
TRACE_START("data_load");
LoadData(...);
TRACE_STOP("data_load");

TRACE_START("compute");
Compute(...);
TRACE_STOP("compute");

TRACE_START("data_store");
StoreData(...);
TRACE_STOP("data_store");

TRACE打点可以分析不同阶段的耗时，找出瓶颈。

6.1.4 内存检查工具

内存泄漏检测

内存泄漏是常见问题，可以用工具检测：

# 概念性命令
# 用Valgrind检测（如果支持）
valgrind --leak-check=full ./your_program

CANN也提供了内存检查工具，可以检测内存泄漏、越界访问等问题。

越界访问检测

越界访问会导致程序崩溃或结果错误，要仔细检查：

边界检查：访问数组或Tensor的时候，要检查索引是否越界。

// 概念性示例
if (index >= 0 && index < size) {
    value.Load(tensor, index);
} else {
    printf("Error: index %d out of range [0, %d)\n", index, size);
}

断言：用assert检查条件，如果条件不满足就报错：

// 概念性示例
#include <cassert>

assert(index >= 0 && index < size);
value.Load(tensor, index);

内存对齐检查

内存不对齐会影响性能，要检查对齐：

// 概念性示例
void CheckAlignment(void* ptr, size_t alignment) {
    if ((uintptr_t)ptr % alignment != 0) {
        printf("Warning: pointer %p is not aligned to %zu bytes\n", ptr, alignment);
    }
}

---

6.2 测试方法

6.2.1 单元测试

什么是单元测试

单元测试就是测试算子的各个功能模块，确保每个模块都正确。

测试范围：单元测试通常测试单个函数或单个功能，不测试整个系统。

测试目标：确保每个功能模块都正确，为后续集成测试打基础。

单元测试设计

设计单元测试要考虑：

测试用例：设计各种测试用例，包括正常情况、边界情况、异常情况。

测试数据：准备测试数据，包括正常数据、边界数据、异常数据。

预期结果：明确每个测试用例的预期结果，用来验证实际结果。

// 概念性示例：Add算子的单元测试
void TestAddKernel() {
    // 准备测试数据
    float input1[] = {1.0f, 2.0f, 3.0f, 4.0f};
    float input2[] = {5.0f, 6.0f, 7.0f, 8.0f};
    float expected[] = {6.0f, 8.0f, 10.0f, 12.0f};
    float output[4];
    
    // 调用算子
    AddKernel(input1, input2, output, 4);
    
    // 验证结果
    for (int i = 0; i < 4; i++) {
        assert(abs(output[i] - expected[i]) < 1e-5);
    }
    
    printf("AddKernel test passed\n");
}

测试框架

可以用测试框架组织测试：

Google Test：C++的测试框架，功能强大。

自定义框架：也可以自己写简单的测试框架，够用就行。

6.2.2 功能测试

功能测试的目标

功能测试就是测试算子的功能是否正确，输入输出是否符合预期。

测试内容：测试算子的基本功能、各种输入情况、边界情况等。

测试方法：用参考实现（比如NumPy）对比结果，或者用已知的正确答案验证。

功能测试设计

设计功能测试：

正常输入测试：测试正常的输入，验证基本功能。

边界输入测试：测试边界值，比如空输入、最大值、最小值等。

异常输入测试：测试异常输入，比如NaN、Inf、负数等，看算子怎么处理。

多维度测试：测试不同形状、不同大小的输入，确保算子在各种情况下都正确。

// 概念性示例：功能测试
void FunctionalTest() {
    // 测试1：正常输入
    TestNormalInput();
    
    // 测试2：边界输入
    TestBoundaryInput();
    
    // 测试3：异常输入
    TestExceptionInput();
    
    // 测试4：不同形状
    TestDifferentShapes();
}

6.2.3 性能测试

性能测试的目标

性能测试就是测试算子的性能，看执行时间、内存使用、NPU利用率等指标。

性能指标：执行时间、吞吐量、延迟、内存带宽利用率、NPU利用率等。

性能基准：和参考实现对比，或者和标准算子对比，看性能是否达标。

性能测试方法

性能测试的方法：

多次运行：多次运行算子，取平均值，减少误差。

预热：先运行几次预热，让系统稳定，再开始测试。

统计信息：收集统计信息，比如最大最小时间、平均值、标准差等。

// 概念性示例：性能测试
void PerformanceTest() {
    const int num_runs = 100;
    double times[num_runs];
    
    // 预热
    for (int i = 0; i < 10; i++) {
        RunKernel();
    }
    
    // 测试
    for (int i = 0; i < num_runs; i++) {
        double start = GetTime();
        RunKernel();
        double end = GetTime();
        times[i] = end - start;
    }
    
    // 统计
    double avg_time = CalculateAverage(times, num_runs);
    double min_time = CalculateMin(times, num_runs);
    double max_time = CalculateMax(times, num_runs);
    
    printf("Average time: %f ms\n", avg_time);
    printf("Min time: %f ms\n", min_time);
    printf("Max time: %f ms\n", max_time);
}

6.2.4 精度测试

精度测试的目标

精度测试就是测试算子的数值精度，看结果是否足够准确。

精度指标：绝对误差、相对误差、最大误差、平均误差等。

精度要求：不同的应用对精度要求不同，要明确精度要求。

精度测试方法

精度测试的方法：

参考对比：和参考实现（比如FP32版本）对比，计算误差。

已知答案：用已知正确答案的数据测试，验证精度。

误差分析：分析误差分布，找出误差大的情况。

// 概念性示例：精度测试
void AccuracyTest() {
    // 准备测试数据
    float input1[], input2[];
    float output_fp16[], output_fp32[];
    
    // FP16版本
    AddKernelFP16(input1, input2, output_fp16, n);
    
    // FP32参考版本
    AddKernelFP32(input1, input2, output_fp32, n);
    
    // 计算误差
    double max_error = 0.0;
    double avg_error = 0.0;
    for (int i = 0; i < n; i++) {
        double error = abs(output_fp16[i] - output_fp32[i]);
        max_error = max(max_error, error);
        avg_error += error;
    }
    avg_error /= n;
    
    printf("Max error: %e\n", max_error);
    printf("Avg error: %e\n", avg_error);
}

6.2.5 边界条件测试

边界条件

边界条件就是输入数据的边界情况，比如空输入、单元素、最大值、最小值等。

空输入：输入为空，看算子怎么处理。

单元素：只有一个元素，看算子是否正确。

最大值最小值：输入最大值最小值，看会不会溢出下溢。

形状边界：不同形状的输入，比如1x1、1xN、Nx1等。

边界测试设计

设计边界测试：

覆盖边界：覆盖各种边界情况，确保都能正确处理。

异常处理：测试异常情况，看算子是否有合适的错误处理。

稳定性：测试边界情况的稳定性，看会不会崩溃或产生错误结果。

// 概念性示例：边界测试
void BoundaryTest() {
    // 测试1：空输入
    TestEmptyInput();
    
    // 测试2：单元素
    TestSingleElement();
    
    // 测试3：最大值
    TestMaxValue();
    
    // 测试4：最小值
    TestMinValue();
    
    // 测试5：不同形状
    TestDifferentShapes();
}

---

6.3 问题排查

6.3.1 常见错误类型

编译错误

编译错误就是代码编译不过，通常是语法错误、类型错误等。

语法错误：括号不匹配、分号缺失等，编译器会报错。

类型错误：类型不匹配、未定义的类型等。

链接错误：找不到函数定义、库文件缺失等。

解决方法：仔细看错误信息，找到问题位置，修复错误。

运行时错误

运行时错误就是程序运行的时候出错，比如段错误、除零错误等。

段错误：访问非法内存，通常是越界访问或空指针。

除零错误：除以零，要检查除数。

内存错误：内存泄漏、重复释放等。

解决方法：用调试工具定位问题，检查内存访问、边界条件等。

逻辑错误

逻辑错误就是程序能运行，但结果不对。

算法错误：算法实现有误，结果不正确。

数据错误：数据处理有误，比如索引计算错误。

精度错误：精度损失太大，结果不准确。

解决方法：用测试用例验证，用日志追踪数据流，找出逻辑错误。

6.3.2 错误定位方法

二分法定位

如果不知道问题在哪，可以用二分法，逐步缩小范围：

注释代码：注释掉一部分代码，看问题是否消失，缩小问题范围。

添加日志：在关键位置添加日志，看程序执行到哪里了。

简化输入：用简单的输入测试，看问题是否复现。

日志追踪

用日志追踪程序执行和数据流：

执行流程：在关键位置打印日志，看程序执行流程。

数据追踪：打印关键变量的值，看数据是否正确。

条件追踪：在条件判断处打印日志，看条件是否满足。

// 概念性示例：日志追踪
void DebugKernel(...) {
    printf("=== Kernel Start ===\n");
    
    printf("Input size: %d\n", input_size);
    for (int i = 0; i < min(10, input_size); i++) {
        printf("Input[%d] = %f\n", i, input[i]);
    }
    
    // 计算
    Compute(...);
    
    printf("Output size: %d\n", output_size);
    for (int i = 0; i < min(10, output_size); i++) {
        printf("Output[%d] = %f\n", i, output[i]);
    }
    
    printf("=== Kernel End ===\n");
}

对比测试

对比测试就是和参考实现对比，找出差异：

参考实现：用已知正确的实现（比如NumPy）作为参考。

逐元素对比：对比每个输出元素，找出差异。

差异分析：分析差异的原因，找出问题。

6.3.3 性能瓶颈分析

识别瓶颈

性能瓶颈就是影响性能的关键部分，要识别出来：

时间分析：用性能分析工具，看哪个部分耗时最多。

资源分析：看NPU利用率、内存带宽利用率等，找出资源瓶颈。

热点分析：找出代码的热点，通常是循环或频繁调用的函数。

瓶颈类型

常见的性能瓶颈：

计算瓶颈：计算太慢，需要优化算法或使用更高效的API。

内存瓶颈：内存访问太慢，需要优化内存访问模式。

同步瓶颈：同步等待太多，需要减少同步或优化同步方式。

负载不均衡：负载不均衡，需要优化负载分配。

优化策略

针对不同的瓶颈，采用不同的优化策略：

计算瓶颈：优化算法、使用融合指令、向量化等。

内存瓶颈：优化访问模式、数据重用、预取等。

同步瓶颈：减少同步、异步执行、流水线化等。

负载不均衡：动态分配、工作窃取、自适应分块等。

6.3.4 优化建议

优化原则

优化的时候要遵循一些原则：

先正确后优化：先保证功能正确，再优化性能。

测量驱动：用数据说话，测量优化效果，不要凭感觉。

逐步优化：一次优化一点，逐步改进，不要一次改太多。

保持可读性：优化不要牺牲代码可读性，要能维护。

优化流程

优化的流程：

1. 性能分析：用工具分析性能，找出瓶颈。

2. 制定方案：根据瓶颈，制定优化方案。

3. 实施优化：实施优化，修改代码。

4. 验证效果：测试优化效果，看性能是否提升。

5. 迭代优化：如果效果不够，继续优化。

常见优化技巧

一些常见的优化技巧：

向量化：尽量用向量化API，一次处理多个元素。

融合操作：用融合指令，减少指令数和内存访问。

数据重用：尽量重用数据，减少重复加载。

预取：提前加载数据，让加载和计算重叠。

分块：合理分块，充分利用Local Memory。

并行：充分利用多核，并行处理。

---

学习检查点

学完这一篇，你应该能做到这些：

掌握调试技术，包括日志输出、断点调试、性能分析工具、内存检查工具的使用。理解测试方法，能够设计单元测试、功能测试、性能测试、精度测试、边界条件测试。掌握问题排查方法，能够定位常见错误、分析性能瓶颈、制定优化方案。能够对算子进行全面的测试和调试，确保功能正确和性能达标。

实践练习

调试实践：在一个有bug的算子中，用日志和调试工具定位问题。练习使用msprof工具分析性能，找出瓶颈。

测试实践：为一个算子设计完整的测试用例，包括单元测试、功能测试、性能测试、精度测试、边界条件测试。执行测试，分析结果，修复发现的问题。

问题排查实践：模拟一些常见错误，练习定位和修复。分析一个性能不佳的算子，找出瓶颈并优化。

综合实践：开发一个完整的算子，从编写到测试到优化，走完整个流程。总结调试和测试的经验，形成自己的方法论。

---

下一步：掌握了调试和测试后，就可以开始实战项目了。下一章会讲如何实现完整的算子，从简单到复杂，从基础到优化，到时候你就能独立开发高质量的算子了。

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

报名链接:https://www.hiascend.com/developer/activities/cann20252

社区地址：https://www.hiascend.com/developer