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: 102402运行验证时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]昇腾文档中心：昇腾社区-官网丨昇腾万里 让智能无所不及
[2]昇腾社区在线课程：开发者主页-昇腾社区
[3]昇腾论坛：https://www.hiascend.com/forum
