printf
产品支持情况
产品 |
是否支持 |
|---|---|
Atlas 350 加速卡 |
√ |
√ |
|
√ |
|
x |
|
x |
|
x |
|
x |
功能说明
本接口提供SIMD和SIMT VF（SIMT VF仅支持Atlas 350 加速卡）调试场景下的格式化输出功能。
在算子Kernel侧的实现代码中，需要输出日志信息时，调用printf接口打印相关内容。
函数原型
1 2 | template <class... Args> __aicore__ inline void printf( const __gm__char* fmt, Args&&... args) |
参数说明
返回值说明
无
约束说明
- 本接口不支持打印除换行符之外的其他转义字符。
- SIMD场景下，单次调用本接口打印的数据总量不可超过打印大小限制，默认为32KB。使用时应注意，如果超出这个限制，则数据不会被打印。您可以通过acl.json中的"simd_printf_fifo_size_per_core"字段进行配置，配置范围最小为1KB，最大为64MB。当打印数据量较大时，建议增加缓存空间。
- SIMT VF中printf功能需要占用额外的Global Memory空间用于数据缓存，缓存空间大小默认为2MB。您可以通过acl.json中的"simt_printf_fifo_size"字段进行配置，配置范围最小为1MB，最大为64MB。当打印数据量较大时，建议增加缓存空间。
- printf在SIMT VF中调用，会占用SIMT栈空间，请合理控制调用printf次数，防止SIMT栈溢出。
- 在仿真环境下，使用printf接口会增加算子运行时间，通过在VF代码中判断线程ID，可以仅在部分线程中打印调试信息，减少重复内容的打印，更有利于调试。
需要包含的头文件
使用该接口需要包含"utils/debug/asc_printf.h"头文件。
1 | #include "utils/debug/asc_printf.h" |
SIMD示例
1 2 3 4 5 6 7 8 9 10 11 12 13 | #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模式下，程序运行时打印效果如下：
1 2 3 4 5 6 7 8 9 10 11 12 13 | 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 ...... |
SIMT示例
1 2 3 4 5 6 7 8 9 10 11 12 | #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模式下，程序运行时打印效果如下：
1 2 3 4 5 6 7 8 9 10 11 | 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 ...... |
父主题： 调测接口