CPU域调试

CPU侧验证核函数

在算子部署在NPU上之前，用户可以在CPU上进行功能的基本验证。下面代码以add_custom算子为例，介绍算子核函数在CPU侧验证时，算子调用的应用程序如何编写。您在实现自己的应用程序时，需要关注由于算子核函数不同带来的修改，包括算子核函数名，入参出参的不同等，合理安排相应的内存分配、内存拷贝和文件读写等，相关API的调用方式直接复用即可。

1. 按需包含头文件，通过ASCENDC_CPU_DEBUG宏区分CPU和NPU侧需要包含的头文件。

   1 2 3 4 5 6 7

   #include "data_utils.h" #ifndef ASCENDC_CPU_DEBUG #include "acl/acl.h" #else #include "tikicpulib.h" extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z); // 核函数声明 #endif

2. CPU侧运行验证。完成算子核函数CPU侧运行验证的步骤如下：
   图1 CPU侧运行验证步骤
   

   GmAlloc、ICPU_RUN_KF、GmFree等接口说明请参考调测接口。

   1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25

   int32_t main(int32_t argc, char* argv[]) { uint32_t blockDim = 8; size_t inputByteSize = 8 * 2048 * sizeof(uint16_t); size_t outputByteSize = 8 * 2048 * sizeof(uint16_t); // 使用GmAlloc分配共享内存，并进行数据初始化 uint8_t* x = (uint8_t*)AscendC::GmAlloc(inputByteSize); uint8_t* y = (uint8_t*)AscendC::GmAlloc(inputByteSize); uint8_t* z = (uint8_t*)AscendC::GmAlloc(outputByteSize); ReadFile("./input/input_x.bin", inputByteSize, x, inputByteSize); ReadFile("./input/input_y.bin", inputByteSize, y, inputByteSize); // 矢量算子需要设置内核模式为AIV模式 AscendC::SetKernelMode(KernelMode::AIV_MODE); // 调用ICPU_RUN_KF调测宏，完成核函数CPU侧的调用 ICPU_RUN_KF(add_custom, blockDim, x, y, z); // 输出数据写出 WriteFile("./output/output_z.bin", z, outputByteSize); // 调用GmFree释放申请的资源 AscendC::GmFree((void *)x); AscendC::GmFree((void *)y); AscendC::GmFree((void *)z); return 0; }

gdb调试

可使用gdb单步调试算子计算精度。由于cpu调测已转为多进程调试，每个核都会拉起独立的子进程，故gdb需要转换成子进程调试的方式。针对 Atlas 推理系列产品 、 Atlas 训练系列产品 ，每个核会拉起1个子进程。针对 Atlas A2 训练系列产品 / Atlas 800I A2 推理产品 /A200I A2 Box 异构组件，每个核会拉起3个子进程，1个Cube，2个Vector。

• 调试单独一个子进程

  启动gdb，示例中的add_custom_cpu为CPU域的算子可执行文件，参考修改并执行一键式编译运行脚本，将一键式编译运行脚本中的run-mode设置成cpu，即可编译生成CPU域的算子可执行文件。

  gdb启动后，首先设置跟踪子进程，之后再打断点，就会停留在子进程中，但是这种方式只会停留在遇到断点的第一个子进程中，其余子进程和主进程会继续执行直到退出。涉及到核间同步的算子无法使用这种方法进行调试。

  gdb --agrs add_custom_cpu  // 启动gdb，add_custom_cpu为算子可执行文件
  (gdb) set follow-fork-mode child

• 调试多个子进程

  如果涉及到核间同步，那么需要能同时调试多个子进程。

  在gdb启动后，首先设置调试模式为只调试一个进程，挂起其他进程。设置的命令如下：

  1	(gdb) set detach-on-fork off

  查看当前调试模式的命令为：

  1	(gdb) show detach-on-fork

  中断gdb程序要使用捕捉事件的方式，即gdb程序捕捉fork这一事件并中断。这样在每一次起子进程时就可以中断gdb程序。设置的命令为：

  1	(gdb) catch fork

  当执行r后，可以查看当前的进程信息：

  1 2 3

  (gdb) info inferiors Num Description * 1 process 19613

  可以看到，当第一次执行fork的时候，程序断在了主进程fork的位置，子进程还未生成。

  执行c后，再次查看info inferiors，可以看到此时第一个子进程已经启动。

  1 2 3 4

  (gdb) info inferiors Num Description * 1 process 19613 2 process 19626

  这个时候可以使用切换到第二个进程，也就是第一个子进程，再打上断点进行调试，此时主进程是暂停状态：

  1 2 3 4 5 6

  (gdb) inferior 2 [Switching to inferior 2 [process 19626] ($HOME/demo)] (gdb) info inferiors Num Description 1 process 19613 * 2 process 19626

  请注意，inferior后跟的数字是进程的序号，而不是进程号。

  如果遇到同步阻塞，可以切换回主进程继续生成子进程，然后再切换到新的子进程进行调试，等到同步条件完成后，再切回第一个子进程继续执行。

如下是调试一个单独子进程的命令样例：

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

使用printf打印命令打印

1
2

printf("xLocal size: %d\n", xLocal.GetSize()); 
printf("tileLength: %d\n", tileLength);