01核函数运行验证时算子存在精度问题

现象描述

在进行算子NPU域的运行验证时，通过md5sum等方式进行算子精度比对，实际数据和真值数据不一致，算子存在精度问题。本示例中通过md5sum来进行精度比对，打印出的真值数据和实际输出数据的md5值不一致，具体打印信息如下：

md5sum:
45e17ee4c068a655be2af4d8c3a1f191  output/golden.bin
6a99e41a84b14dd04f32730ceb9a3988  output/output_y.bin

可能原因

算子出现精度问题，一般是由于算子的实现逻辑有误。

处理步骤

Ascend C提供孪生调试的功能，通过CPU域的功能验证、gdb单步调试、printf数值打印来定位算子的实现逻辑问题。本样例仅展示了可能会出现的场景，便于演示定位步骤。实际使用过程中，请根据代码情况进行调试。

1. 进行CPU域的功能验证，观察是否有日志报错。

编写CPU侧的运行验证代码，并进行运行验证。得到CPU域的精度比对结果如下：

md5sum:
45e17ee4c068a655be2af4d8c3a1f191 output/golden.bin
5d6e1aec686b28bd3839dbcd5caaa8b2 output/output_y.bin

可以看出CPU域的精度比对也存在不一致的问题，然后观察是否有打屏日志报错，可搜索关键词"failed"。比如，下图的报错示例指示，错误出现在代码中调用LeakyRelu接口的地方。

leakyrelu_custom_cpu: /usr/local/Ascend/CANN-7.0/x86_64-linux/tikcpp/tikcfw/interface/kernel_operator_vec_binary_scalar_intf.h:447: void AscendC::LeakyRelu(const AscendC::LocalTensor<T>&, const AscendC::LocalTensor<T>&, const T&, const int32_t&) [with T = float16::Fp16T; int32_t = int]: Assertion `false && "check vlrelu instr failed"' failed

通过上述报错日志，一般只能定位到报错的代码行，无法明确具体错误，接下来需要通过gdb调试的方式或者printf打印的方式进一步精确定位。

2. gdb调试。下面的样例展示了拉起leakyrelu算子CPU侧运行程序的样例，该样例程序会直接抛出异常，直接gdb运行，查看调用栈信息分析定位即可。其他场景下您可以使用gdb打断点等基本操作进行调试。

1) 使用gdb拉起待调试程序，进入gdb界面进行debug。

gdb leakyrelu_custom_cpu

2) 单独调试一个子进程。

(gdb) set follow-fork-mode child

3) 运行程序。

(gdb) r

4) 通过bt查看程序调用栈。

(gdb) bt

5) 查看具体层的堆栈信息，打印具体变量的值。本示例中，打印了tileLength为1024，该程序中表示需要处理1024个half类型的数，大小为1024*sizeof(half)=2048字节；输入Tensor xLocal的值，其中dataLen表示LocalTensor的size大小为1024字节，只能计算1024字节的数据。可以看出两者的长度不匹配，由此可以定位问题。

(gdb) f 5
#5 0x000055555555d364 in KernelLeakyRelu::Compute (this=0x7fffffffd7d0, progress=0) at /root/AscendC_DemoCode-master/precision-error/vector/leakyrelu_custom.cpp:59
59 LeakyRelu(yLocal, xLocal, scalar, tileLength);
(gdb) p tileLength
$1 = 1024
(gdb) p xLocal
$1 = {<AscendC::BaseTensor<float16::Fp16T>> = {<No data fields>}, address_ = {logicPos = 9 '\t', bufferHandle = 0x7fffffffd930 "\003\005\377\377", dataLen = 1024,bufferAddr = 0,absAddr = ...}

3. printf打印。在合适的位置增加变量打印。样例代码如下：

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

可以看到有如下打屏日志输出，打印了tileLength为1024，该程序中表示需要处理1024个half类型的数；输入Tensor xLocal的size大小，为512，表示只能计算512个half类型的数。可以看出两者的长度不匹配，由此可以定位问题。

xLocal size: 512
tileLength: 1024

02运行验证时AllocTensor/FreeTensor失败

现象描述

通过NPU进行核函数的运行验证时，出现挂死现象；通过CPU进行核函数的运行验证时，出现AllocTensor/FreeTensor失败的报错，日志报错和调用栈打印如下：

[ERROR][Core_0][/usr/local/Ascend/latest/x86_64-linux/tikcpp/tikcfw/interface/kernel_tpipe.h:730][AllocEventID][321678] current size is 4, max buffer number in same queue position is 4
[ERROR][CORE_0][pid 321674] error happened! =========
SIGABRT Signal (Abort Signal from abort) catched, backtrace info:
[#0] 0x000000000001e7c0: handler(int) at /usr/local/Ascend/latest/tools/tikicpulib/lib/include/kern_fwk.h:105
[#1] 0x0000000000017c4f: signed char AscendC::TPipe::AllocEventID<(AscendC::HardEvent)5>() at /usr/local/Ascend/latest/x86_64-linux/tikcpp/tikcfw/interface/kernel_tpipe.h:733
[#2] 0x000000000001426d: AscendC::TQueBind<(AscendC::TPosition)0, (AscendC::TPosition)9, 4, 0>::FreeBuffer(unsigned char*) at /usr/local/Ascend/latest/x86_64-linux/tikcpp/tikcfw/interface/kernel_tpipe.h:1217
[#3] 0x0000000000011058: void AscendC::TQueBind<(AscendC::TPosition)0, (AscendC::TPosition)9, 4, 0>::FreeTensor<float16::Fp16T>(AscendC::LocalTensor<float16::Fp16T>&) at /usr/local/Ascend/latest/x86_64-linux/tikcpp/tikcfw/interface/kernel_tpipe.h:1237
[#4] 0x000000000000dfde: KernelAdd::Compute(int) at /home/xxxx/xxxx.cpp:59
[#5] 0x000000000000dd1c: KernelAdd::Process() at /home/xxxx/xxxx.cpp:37 (discriminator 2)
...

可能原因

根据日志信息“current size is 4, max buffer number in same queue position is 4”可以明确该问题是因为同一个TPosition上QUE Buffer的数量超出限制导致。

同一个TPosition上QUE Buffer的数量根据AI处理器型号的不同，有数量约束。申请Buffer时，需要满足该约束。

Atlas 训练系列产品、Atlas推理系列产品（Ascend 310P处理器）AI Core不超过4块。

Atlas A2训练系列产品/Atlas 300I A2推理产品不超过8块。

不满足该约束，可能会在后续使用AllocTensor/FreeTensor可能会出现分配资源失败。比如：

TQue<TPosition::VECIN, 1> que0;
TQue<TPosition::VECIN, 1> que1;
// 不建议：
// 比如，算子有6个输入，需要申请6块buffer
// 通过2个队列为其申请内存，分别为que0、que1分配3块,申请VECIN position上的buffer总数为6
// 针对Atlas 训练系列产品、Atlas推理系列产品（Ascend 310P处理器）AI Core同一个TPosition上QUE Buffer的数量限制为4，超出该限制，在后续使用AllocTensor/FreeTensor可能会出现分配资源失败。
pipe.InitBuffer(que0, 3, len);
pipe.InitBuffer(que1, 3, len);

处理步骤

如果确实有多块buffer使用, 可以将多个buffer合并到一块buffer, 通过偏移使用。样例如下：

// 此时建议通过以下方法解决：
// 如果确实有多块buffer使用, 可以将多个buffer合并到一块buffer, 通过偏移使用
pipe.InitBuffer(que0, 1, len * 3)
pipe.Initbuffer(que1, 1, len * 3)
/*
 * 分配出3块内存大小的LocalTensor, local1的地址为que0中buffer的起始地址，
 * local2的地址为local1的地址偏移len后的地址，local3的地址为local1的地址偏移
 * len * 2的地址
 */
int32_t offset1 = len;
int32_t offset2 = len * 2;
LocalTensor<T> local1 = que0.AllocTensor<T>();
LocalTensor<T> local2 = local1[offset1];
LocalTensor<T> local3 = local1[offset2];

03 kernel侧获取Tiling信息不正确

现象描述

通过算子在kernel侧实现代码中添加PRINTF打印发现kernel侧获取的Tiling信息不正确。

比如如下样例，增加的打印代码如下：

PRINTF("tiling_data.totalLength: %d tiling_data.tileNum: %d.\n",tiling_data.totalLength, tiling_data.tileNum);

打印的Tiling数据如下，全为0：

tiling_data.totalLength: 0 tiling_data.tileNum: 0.

可能原因

kernel侧获取Tiling信息不正确的原因一般有以下两种：

• host侧计算Tiling的逻辑不正确

• kernel侧核函数的参数未按照正确顺序填写

处理步骤

1. 参考如下示例，打印TilingData的数据，确认host侧序列化保存的TilingData是否正确。如果此时打印值有误，说明Tiling的计算逻辑可能不正确，需要进一步检查host侧Tiling实现代码，排查计算逻辑是否有误。

std::out<<*reinterpret_cast<uint32_t *>(context->GetRawTilingData()->GetData())<<std::endl; //按照实际数据类型打印TilingData第一个参数值，如需确认其他值，取值指针向后偏移即可

2. 如果上一步骤中打印的TilingData正确，需要排查kernel侧核函数的参数是否按照正确顺序填写。

使用msopgen工具创建算子工程，并基于工程进行kernel侧算子开发时，核函数的定义模板已通过msopgen工具自动生成，样例如下所示。参数按照“输入、输出、workspace、tiling”的顺序排布。请检查是否调整过参数顺序导致和正确顺序不一致。

#include "kernel_operator.h"
extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling) {
GET_TILING_DATA(tiling_data, tiling);// 获取Tiling参数
// TODO: user kernel impl
}

04 更多介绍

[1]昇腾文档中心：https://www.hiascend.com/zh/document

[2]昇腾社区在线课程：https://www.hiascend.com/zh/edu/courses

[3]昇腾论坛：https://www.hiascend.com/forum