cann/asc-devkit printf调试接口

printf【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit产品支持情况Ascend 950PR/Ascend 950DT支持Atlas A3 训练系列产品/Atlas A3 推理系列产品支持Atlas A2 训练系列产品/Atlas A2 推理系列产品支持Atlas 200I/500 A2 推理产品不支持Atlas 推理系列产品AI Core不支持Atlas 推理系列产品Vector Core不支持Atlas 训练系列产品不支持Kirin X90不支持Kirin 9030不支持功能说明本接口提供SIMD和SIMT编程调试场景下的格式化输出功能。在算子Kernel侧的实现代码中需要输出日志信息时调用printf接口打印相关内容。[!CAUTION]注意 printfPRINTF接口打印功能会对算子实际运行的性能带来一定影响通常在调测阶段使用。开发者可以按需通过设置ASCENDC_DUMP0的方式关闭打印功能。函数原型template class... Args __aicore__ inline void printf( const __gm__ char* fmt, Args... args)以下接口为simd_vf中所使用的printf接口仅支持Ascend 950PR/Ascend 950DT。template class... Args __simd_callee__ inline void printf( const __ubuf__ char* fmt, Args... args)参数说明参数名输入/输出描述fmt输入格式控制字符串包含两种类型的对象普通字符和转换说明。普通字符将原样不动地打印输出。转换说明并不直接输出而是用于控制printf中参数的转换和打印。每个转换说明都由一个百分号字符%开始以转换说明结束从而说明输出数据的类型。支持的转换类型包括%d、%ld、%lld、%i、%li、%lli输出十进制数支持打印的数据类型int8_t、int16_t、int32_t、int64_t%f、%F输出浮点数支持打印的数据类型float、half、bfloat16_t%x、%lx、%llx输出十六进制整数支持打印的数据类型int8_t、int16_t、int32_t、int64_t、uint8_t、uint16_t、uint32_t、uint64_t%s输出字符串%u、%lu、%llu输出unsigned类型数据支持打印的数据类型uint8_t、uint16_t、uint32_t、uint64_t%p输出指针地址注意上文列出的数据类型是NPU域调试支持的数据类型CPU域调试时支持的数据类型和C/C规范保持一致。args输入附加参数个数和类型可变的参数列表根据不同的fmt字符串函数可能需要一系列的附加参数每个参数包含了一个要被插入的值替换了fmt参数中指定的每个%标签。参数的个数应与%标签的个数相同。返回值说明无约束说明本接口不支持打印除换行符之外的其他转义字符。SIMT中printf功能需要占用额外的Global Memory空间用于数据缓存缓存空间大小默认为2MB。您可以通过acl.json中的simt_printf_fifo_size字段进行配置配置范围最小为1MB最大为64MB。当打印数据量较大时建议增加缓存空间。在SIMT中调用printf接口会占用SIMT栈空间请控制每次printf打印的数据量防止SIMT栈溢出。使用printf接口会增加算子运行时间请控制打印数据量避免算子超时。建议在代码中判断线程ID仅在部分线程中打印调试信息减少重复内容的打印更有利于调试。SIMD场景下单次调用本接口打印的数据总量不可超过打印大小限制默认为32KB。使用时应注意如果超出这个限制则数据不会被打印。您可以通过acl.json中的simd_printf_fifo_size_per_core字段进行配置配置范围最小为1KB最大为64MB。当打印数据量较大时建议增加缓存空间。pytorch调用和算子入图场景暂不支持该配置。SIMD场景下根据算子执行方式的不同printf的打印结果输出方式不同。动态图或者单算子直调场景下待输出内容会被解析并打印在屏幕上静态图场景下整图算子需要全下沉到NPU侧执行无法直接调用接口打印出单个算子的信息因此需要在模型执行完毕后将待输出内容落盘在dump文件中dump文件需要通过工具解析为可读内容。dump文件落盘路径按照优先级排列如下如果开启了Data Dump功能dump文件落盘到开发者配置的dump_path路径下。如何开启Dump功能依赖于具体的网络运行方式。以TensorFlow在线推理为例通过enable_dump、dump_path、dump_mode等参数进行配置。配置方式可参考《TensorFlow 2.6.5模型迁移》中的API参考 TF Adapter接口2.x npu.global_options 配置参数说明章节。如果未开启Data Dump功能但配置了ASCEND_WORK_PATH环境变量dump文件落盘到ASCEND_WORK_PATH下的printf目录下。如果未开启Data Dump功能也没有配置ASCEND_WORK_PATH环境变量dump文件落盘到当前程序执行目录下的printf路径下。落盘dump文件需要使用工具解析为用户可读内容使用show_kernel_debug_data工具将dump二进制文件解析为用户可读内容命令格式如下。show_kernel_debug_data bin_file output_dir由于在simd vf中不能直接访问gm地址因此使用simd vf中的printf时需要先定义一个ubuf的字符串该字符串即为原来的fmt变量。__ubuf__ const char* fmt simd vf: int%d, uint%u, float%f, string%s\n;需要包含的头文件使用该接口需要包含utils/debug/asc_printf.h头文件。#include utils/debug/asc_printf.hSIMD调用示例#include utils/debug/asc_printf.h // SIMD printf __global__ __mix__(1, 2) void hello_world() { // print string printf(hello world device\n); // print int printf(fmt string int: %d\n, 0x123); // print float float b 3.14; printf(fmt string float: %f\n, b); }NPU模式下程序运行时打印效果如下hello world device fmt string int: 291 fmt string float: 3.140000 hello world device fmt string int: 291 fmt string float: 3.140000 hello world device fmt string int: 291 fmt string float: 3.140000 hello world device fmt string int: 291 fmt string float: 3.140000 ......SIMD VF调用示例#include kernel_operator.h #include utils/debug/asc_printf.h __simd_vf__ inline void SimdVfPrint() { __ubuf__ const char* fmt simd vf: int%d, uint%u, float%f, string%s\n; printf(fmt, 1, 2U, 5.0f, AscendC); }NPU模式下程序运行时打印效果如下simd vf: int1, uint2, float5.000000, stringAscendC simd vf: int1, uint2, float5.000000, stringAscendC simd vf: int1, uint2, float5.000000, stringAscendC ......SIMT示例#include kernel_operator.h #include simt_api/asc_simt.h #include utils/debug/asc_printf.h // asc_vf_call调用时dim3参数dim3(8, 2, 8) __simt_vf__ __launch_bounds__(128) inline void SimtCompute() { int x threadIdx.x; int y threadIdx.y; int z threadIdx.z; printf(simt vf: d: (%d, %d, %d), f: %f, s: %s\n, x, y, z, 3.14f, pass); }NPU模式下程序运行时打印效果如下simt vf: d: (0, 0, 0), f: 3.140000, s: pass simt vf: d: (0, 0, 1), f: 3.140000, s: pass simt vf: d: (0, 0, 2), f: 3.140000, s: pass simt vf: d: (0, 0, 3), f: 3.140000, s: pass simt vf: d: (0, 0, 4), f: 3.140000, s: pass simt vf: d: (0, 0, 5), f: 3.140000, s: pass simt vf: d: (0, 0, 6), f: 3.140000, s: pass simt vf: d: (0, 0, 7), f: 3.140000, s: pass simt vf: d: (0, 1, 0), f: 3.140000, s: pass simt vf: d: (0, 1, 1), f: 3.140000, s: pass ......【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考