3天上手Ascend C编程 | Day3 Ascend C算子调试调优方法

举报
昇腾CANN 发表于 2023/07/28 14:54:23 2023/07/28
【摘要】 Ascend C提供孪生调试方法，即在cpu侧创建一个npu的模型并模拟它的计算行为，用来进行业务功能调试。相同的算子代码可以在cpu域调试精度，npu域调试性能。

本文分享自《 【2023 · CANN训练营第一季】——Ascend C算子开发入门——第三次课》，作者：dayao

Ascend C是CANN针对算子开发场景推出的编程语言，原生支持C和C++标准规范，最大化匹配用户开发习惯；通过多层接口抽象、自动并行计算、孪生调试等关键技术，极大提高算子开发效率，助力AI开发者低成本完成算子开发和模型调优部署。

时间充足的小伙伴推荐去看Ascend C官方教程：

想省时省力快速入门可以看这篇文章，为你系统化梳理AscendC编程最重要的知识点，3天快速上手不迷路。
第3天的学习内容要点如下：

固定shape算子改写成动态shape算子


在调试之前，我们先来讲下如何将第2次课固定shape的add算子改造成动态shape的add算子，以符合实际算子运行场景。

1、什么是动态shape

动态shape和固定shape是一对概念。固定shape是指，在编译时指定shape大小，运行时不需要指定shape大小。在开发固定shape算子过程中，一个算子源代码可以支持多个固定shape，但需要在编译时明确了shape的实际值。而动态shape则是在编译时不指定shape大小，在运行时传入实际的shape大小，算子编译后的二进制文件支持任意shape，或者是一个或多个shape范围。

2、如何实现动态shape

课程的第2次课，讲述是采用固定shape的加法实例，本次课讲述，如何将固定shape改为动态shape的算子。也就是，将控制形状的BLOCK DIM，TOTAL LENGTH，TILE NUM这些变量做成tiling结构体，作为参数传给核函数。如下所示：

3、动态shape算子的tiling结构体
1）主要操作流程
2）tiling结构体中的信息
  • BLOCK_DIM：并行计算使用的核数
  • TOTAL_LENGTH：总共需要计算的数据个数
  • TILE_NUM：每个核上计算数据分块的个数 
struct AddCustomTilingData {
    uint32_t blockDim;
    uint32_t totalLength;
    uint32_t tileNum;
};

3）动态shape算子的tiling解析函数

4、固定与动态shape实现对比

1）核函数实现对比
固定shape输入的核函数实现： 
constexpr int32_t BLOCK_DIM = 8;                                      // num of core used
constexpr int32_t TOTAL_LENGTH = 8 * 2048;                            // total length of data
constexpr int32_t BLOCK_LENGTH = TOTAL_LENGTH / BLOCK_DIM;            // length computed of each core
constexpr int32_t TILE_NUM = 8;                                       // split data into 8 tiles for each core
constexpr int32_t BUFFER_NUM = 2;                                     // tensor num for each queue
constexpr int32_t TILE_LENGTH = BLOCK_LENGTH / TILE_NUM / BUFFER_NUM; // each tile length is separated to 2 part, due to double buffer
 
// implementation of kernel function
extern "C" __global__ __aicore__ void add_custom(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z)
{
    KernelAdd op;
    op.Init(x, y, z);
    op.Process();
}

动态shape输入的核函数实现：

constexpr int32_t BUFFER_NUM = 2; // tensor num for each queue
 
// implementation of kernel function
extern "C" __global__ __aicore__ void add_custom(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z, __gm__ uint8_t* tiling)
{
    GET_TILING_DATA(tilingData, tiling);
    KernelAdd op;
    op.Init(x, y, z, tilingData.blockDim, tilingData.totalLength, tilingData.tileNum);
    op.Process();
}
2）kernelAdd类实现对比
3）Init()函数实现对比
固定shape输入的Init()函数实现： 
__aicore__ inline void Init(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z)
{
    // get start index for current core, core parallel
    xGm.SetGlobalBuffer((__gm__ half*)x + block_idx * BLOCK_LENGTH);
    yGm.SetGlobalBuffer((__gm__ half*)y + block_idx * BLOCK_LENGTH);
    zGm.SetGlobalBuffer((__gm__ half*)z + block_idx * BLOCK_LENGTH);
    // pipe alloc memory to queue, the unit is Bytes
    pipe.InitBuffer(inQueueX, BUFFER_NUM, TILE_LENGTH * sizeof(half));
    pipe.InitBuffer(inQueueY, BUFFER_NUM, TILE_LENGTH * sizeof(half));
    pipe.InitBuffer(outQueueZ, BUFFER_NUM, TILE_LENGTH * sizeof(half));
}
动态shape输入的Init()函数实现：
__aicore__ inline void Init(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z,
    uint32_t blockDim, uint32_t totalLength, uint32_t tileNum)
{
    this->blockLength = totalLength / blockDim;
    this->tileNum = tileNum;
    this->tileLength = this->blockLength / tileNum / BUFFER_NUM;
    // get start index for current core, core parallel
    xGm.SetGlobalBuffer((__gm__ half*)x + block_idx * this->blockLength);
    yGm.SetGlobalBuffer((__gm__ half*)y + block_idx * this->blockLength);
    zGm.SetGlobalBuffer((__gm__ half*)z + block_idx * this->blockLength);
    // pipe alloc memory to queue, the unit is Bytes
    pipe.InitBuffer(inQueueX, BUFFER_NUM, this->tileLength * sizeof(half));
    pipe.InitBuffer(inQueueY, BUFFER_NUM, this->tileLength * sizeof(half));
    pipe.InitBuffer(outQueueZ, BUFFER_NUM, this->tileLength * sizeof(half));
}
4）真值生成脚本对比
5）main.cpp对比
main.cpp差异在于动态shape，需要根据输入的tiling计算出实际的数据大小。
固定shape输入的main函数(部分）： 
size_t inputByteSize = 8 * 2048 * sizeof(uint16_t);
size_t outputByteSize = 8 * 2048 * sizeof(uint16_t);
uint32_t blockDim = 8;
动态shape输入的main函数(部分）： 
uint8_t* tiling = (uint8_t*)addcustom::GmAlloc(tilingSize);
ReadFile("./input/tiling.bin", tilingSize, tiling, tilingSize);uint32_t blockDim = (*(const uint32_t *)(tiling));
 
size_t inputByteSize = blockDim * 2048 * sizeof(uint16_t);
size_t outputByteSize = blockDim * 2048 * sizeof(uint16_t);
 
// =========================================
 
aclrtMallocHost((void**)(&tilingHost), tilingSize);
ReadFile("./input/tiling.bin", tilingSize, tilingHost, tilingSize);
uint32_t blockDim = (*(const uint32_t *)(tilingHost));
size_t inputByteSize = blockDim * 2048 * sizeof(uint16_t);
size_t outputByteSize = blockDim * 2048 * sizeof(uint16_t);
综上，固定shape和动态shape代码文件的差别如下：

5、运行结果

CPU模式下的结果：
NPU模式下的结果：

CPU模式下算子调试技术


Ascend C提供孪生调试方法，即在cpu侧创建一个npu的模型并模拟它的计算行为，用来进行业务功能调试。以此进行业务功能的调试。相同的算子代码可以在CPU模式下进行精度调试，然后无缝切换到NPU模式下运行，主要有两种方法：

1、使用GDB进行调试

source /usr/local/Ascend/ascend-toolkit/set_env.sh
gdb --args add_custom_cpu
set follow-fork-mode child
break add_custom.cpp:45
run
list
backtrace
print i
break add_custom.cpp:56
continue
display xLocal
quit

由于cpu调测已转为多进程调试，每个核都是一个独立的子进程，故gdb需要转换成子进程调试的方式。

1）调试单独一个子进程
在gdb启动后，首先设置跟踪子进程，之后再打断点，就会停留在子进程中，设置的命令为： 
set follow-fork-mode child
但是这种方式只会停留在遇到断点的第一个子进程中，其余子进程和主进程会继续执行直到退出。涉及到核间同步的算子无法使用这种方法进行调试。
2）调试多个子进程
如果涉及到核间同步，那么需要能同时调试多个子进程。
在gdb启动后，首先设置调试模式为只调试一个进程，挂起其他进程。设置的命令如下： 
(gdb) set detach-on-fork off
查看当前调试模式的命令为：
(gdb) show detach-on-fork
中断gdb程序的方式要使用捕捉事件的方式，即gdb程序监控fork这一事件并中断。这样在每一次起子进程时就可以中断gdb程序。设置的命令为：(gdb) catch fork
当执行r后，可以查看当前的进程信息：
(gdb) info inferiors  
Num  Description
* 1    process 19613
可以看到，当第一次执行fork的时候，程序断在了主进程fork的位置，子进程还未生成。
执行c后，再次查看info inferiors，可以看到此时第一个子进程已经启动。 
(gdb) info inferiors  
Num  Description 
* 1    process 19613 
2    process 19626
这个时候可以使用切换到第二个进程，也就是第一个子进程，再打上断点进行调试，此时主进程是暂停状态： 
(gdb) inferior 2[Switching to inferior 2 
[process 19626] ($HOME/demo)]
(gdb) info inferiors  Num  Description  
1    process 19613
* 2    process 19626
请注意，inferior后跟的数字是进程的序号，而不是进程号。
如果遇到同步阻塞，可以切换回主进程继续生成子进程，然后再切换到新的子进程进行调试，等到同步条件完成后，再切回第一个子进程继续执行。

2、使用printf或者std::cout

在CPU代码侧直接插入C/C++的打印命令，如printf、std:.cout，但注意NPU模式下不支持打印语句，所以需要添加内置宏__CCE KT TEST__ 予以区分。

性能数据采集与分析


当使用内核调用符时，会生成相应的二进制可执行文件，可以使用性能采集工具运行NPU模式下生成的可执行文件，从而采集Ascend C算子在昇腾平台上的性能数据。​

1.设置环境变量(昇腾实际安装的位置)：

source /usr/local/Ascend/ascend-toolkit/../set_env.sh

2.测试NPU模式下的Ascend C算子，保证其精度正确，生成对应的二进制可执行文件：

bash run.sh add_tik2 ascend910 AiCore npu

3.使用msprof工具采集性能，详细的介绍请参考昇腾社区文档：性能分析工具使用教程

用户可以根据自身的需要灵活组合性能分析指令，例如：

msprof --application="./add_custom_npu" --output="./out" --ai-core=on --aic-metrics="PipeUtilization"

4.在当前目录会生成out文件夹，在device_0/summary/op_summary_0_1.csv能够看到一些具体的数据：

  • mte2类型指令（DDR->AICORE搬运类指令）的cycle数在所有指令的cycle数中的占用比

  • mte3类型指令（AICORE->DDR搬运类指令）的cycle数在所有指令的cycle数中的占用比

  • vector类型指令（向量类运算指令）的cycle数在所有指令的cycle数中的占用比

接下来就可以针对实际情况进行相应优化。

更多学习资源

好啦，本次分享结束啦，Ascend C的学习资源还有很多，想深入学习的同学可以参考Ascend C官网教程。

3天上手Ascend C编程 | Day1 Ascend C基本概念及常用接口
3天上手Ascend C编程 | Day2 通过Ascend C编程范式实现一个算子实例
3天上手Ascend C编程 | Day3 Ascend C算子调试调优方法
【版权声明】本文为华为云社区用户原创内容，转载时必须标注文章的来源（华为云社区）、文章链接、文章作者等基本信息， 否则作者和本社区有权追究责任。如果您发现本社区中有涉嫌抄袭的内容，欢迎发送邮件进行举报，并提供相关证据，一经查实，本社区将立刻删除涉嫌侵权内容，举报邮箱： cloudbbs@huaweicloud.com
人工智能 开发者 昇腾 机器学习
  • 点赞
  • 收藏
  • 关注作者

评论（0）

0/1000
抱歉，系统识别当前为高风险访问，暂不支持该操作

全部回复

上滑加载中

设置昵称

在此一键设置昵称，即可参与社区互动！

*长度不超过10个汉字或20个英文字符，设置后3个月内不可修改。

*长度不超过10个汉字或20个英文字符，设置后3个月内不可修改。

确认 取消

加入云驻计划，成为创作者

  • 华为云周边好礼
  • 免费体验产品
  • 特殊身份标识
  • 线下官方门票
  • 内部专家零距离
  • 与10000+优质创作者共同成长
立即加入