---
title: Ascend C算子开发常见问题案例-官方技术文章-昇腾社区
description: Ascend C是CANN针对算子开发场景推出的编程语言，原生支持C和C++标准规范，最大化匹配用户开发习惯；通过自动并行计算、孪生调试等关键技术，极大提高算子开发效率，助力AI开发者低成本完成算子开发和模型调优部署。 本期分享几个Ascend C算子开发过程中常见问题的典型案例，并给出原因分析及解决方法。
keywords: Ascend,算子开发常见,问题案例,官方技术文章,昇腾社区,核函数运行验,证时算子存在,精度问题
url: https://www.hiascend.com/developer/techArticles/20240106-1?envFlag=1
section: (其他)
---

# Ascend C算子开发常见问题案例-官方技术文章-昇腾社区

URL: https://www.hiascend.com/developer/techArticles/20240106-1?envFlag=1
描述: Ascend C是CANN针对算子开发场景推出的编程语言，原生支持C和C++标准规范，最大化匹配用户开发习惯；通过自动并行计算、孪生调试等关键技术，极大提高算子开发效率，助力AI开发者低成本完成算子开发和模型调优部署。 本期分享几个Ascend C算子开发过程中常见问题的典型案例，并给出原因分析及解决方法。
关键词: Ascend,算子开发常见,问题案例,官方技术文章,昇腾社区,核函数运行验,证时算子存在,精度问题

官方技术文章 [了解详情](https://www.hiascend.com/zh/developer/techArticles)

Ascend C算子开发常见问题案例

Ascend C算子开发常见问题案例

昇腾CANNAscend C

发表于: 2024/01/06

# 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 [了解详情](https://www.hiascend.com/zh/document)[2]昇腾社区在线课程：https://www.hiascend.com/zh/edu/courses [了解详情](https://www.hiascend.com/zh/edu/courses)[3]昇腾论坛：https://www.hiascend.com/forum [了解详情](https://www.hiascend.com/forum)

本页内容