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

2023-12-29 22:42:05

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]昇腾文档中心:昇腾社区-官网丨昇腾万里 让智能无所不及

[2]昇腾社区在线课程:开发者主页-昇腾社区

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

文章来源:https://blog.csdn.net/m0_71340392/article/details/135292557
本文来自互联网用户投稿,该文观点仅代表作者本人,不代表本站立场。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。