在样例工程使用Ascend C算子调测API
AscendC算子调测API是AscendC提供的调试能力,可进行kernel内部的打印、Tensor内容的查看(Dump)。
关于kernel调测api的详细介绍,可参考DumpTensor和printf.
插入调试代码
修改使用该功能的核函数入口和相关调用代码,增加开启调测功能(
#if defined(ENABLE_ASCENDC_DUMP))的编译时代码,具体可参考examples/sdma/main.cpp。在想进行调试的层级,增加调测API调用,如下所示:
// examples/sdma/main.cpp template <typename T> __global__ __aicore__ void allgather_sdma_tensor(GM_ADDR gva, int elem_size, GM_ADDR dump, bool is_put) { AscendC::TPipe pipe; +#if defined(ENABLE_ASCENDC_DUMP) + AscendC::InitDump(false, dump, ALL_DUMPSIZE); +#endif if ASCEND_IS_AIV { + AscendC::printf("my_pe is %d\n", aclshmem_my_pe()); + AscendC::GlobalTensor<T> gmT; + gmT.SetGlobalBuffer((__gm__ T*)gva, elem_size); + AscendC::DumpTensor(gmT, elem_size, 16); ... } }注意:
ALL_DUMPSIZE及aclCheck等宏和接口定义位于文件examples/utils/debug.h中,其中ALL_DUMPSIZE默认为75MB,用户可根据需要进行自定义修改。
编译运行
打开工具的编译开关
-enable_ascendc_dump, 使能AscendC算子调测API编译算子样例。A2/A3 平台:
bash scripts/build.sh -enable_ascendc_dump -examples
Ascend950 平台:
bash scripts/build.sh -soc_type Ascend950 -enable_ascendc_dump -examples
在examples/sdma目录执行demo:
bash run.sh -pes 2 -type int