在样例工程使用Ascend C算子调测API
AscendC算子调测API是AscendC提供的调试能力,可进行kernel内部的打印、Tensor内容的查看(Dump)。
关于kernel调测api的详细介绍,可参考DumpTensor和printf.
插入调试代码
修改使用该功能的核函数入口和相关调用代码,增加开启调测功能(
#if defined(ENABLE_ASCENDC_DUMP))的编译时代码,具体可参考examples/allgather_matmul/main.cpp。在想进行调试的层级,增加调测API调用,如下所示:
// examples/allgather_matmul/main.cpp #if defined(ENABLE_ASCENDC_DUMP) __global__ __aicore__ void ShmemAllGatherMatmul( uint64_t fftsAddr, GM_ADDR gmA, GM_ADDR gmB, GM_ADDR gmC, GM_ADDR gmSymmetric, uint32_t m, uint32_t n, uint32_t k, GM_ADDR dump) { AscendC::InitDump(false, dump, ALL_DUMPSIZE); #else __global__ __aicore__ void ShmemAllGatherMatmul( uint64_t fftsAddr, GM_ADDR gmA, GM_ADDR gmB, GM_ADDR gmC, GM_ADDR gmSymmetric, uint32_t m, uint32_t n, uint32_t k) { #endif // Set FFTS address AscendC::SetSyncBaseAddr(fftsAddr); + AscendC::printf("fftsAddr is %d\n", fftsAddr); + AscendC::GlobalTensor<ElementB> gmT; + gmT.SetGlobalBuffer((__gm__ ElementB*)gmB, k * n); + AscendC::DumpTensor(gmT, n, 16); ... }
注意:
ALL_DUMPSIZE及aclCheck等宏和接口定义位于文件examples\utils\debug.h中,其中ALL_DUMPSIZE默认为75MB,用户可根据需要进行自定义修改。
编译运行
打开工具的编译开关
-enable_ascendc_dump, 使能AscendC算子调测API编译算子样例。bash scripts/build.sh -enable_ascendc_dump -examples
在examples/allgather_matmul目录执行demo:
bash scripts/run.sh 6,7