printf

Applicability

Function Usage

Provides the formatted output function in SIMD and SIMT VF debugging scenarios (SIMT VF supported only on Atlas 350 Accelerator Card).

In the operator kernel implementation code, call the printf API to print related content when log information needs to be output.

Prototype

1
2

template <class... Args>
__aicore__ inline void printf( const __gm__char* fmt, Args&&... args)

Parameters

Returns

None

Restrictions

• This API does not allow the printing of escape characters, except for newline characters.
• In the SIMD scenario, the total amount of data printed by a single call to this API cannot exceed the printing size limit, which is 32 KB by default. If the limit is exceeded, the data will not be printed. You can configure this limit using the simd_printf_fifo_size_per_core field in the acl.json file. The value ranges from 1 KB to 64 MB. If a large amount of data needs to be printed, you are advised to increase the cache space.
• The printf function in SIMT VF requires additional global memory space for data caching. The default cache space is 2 MB. You can configure the cache space using the simt_printf_fifo_size field in the acl.json file. The value ranges from 1 MB to 64 MB. If a large amount of data needs to be printed, you are advised to increase the cache space.
• Calling printf in SIMT VF consumes SIMT stack space. Control the number of calls to prevent stack overflow.
• In the simulation environment, using the printf API increases operator runtime. By checking thread IDs in VF code, debug information can be printed from selected threads only, reducing redundant output and improving debugging efficiency.

Header Files to Be Included

To use this API, the utils/debug/asc_printf.h header file must be included.

1

#include "utils/debug/asc_printf.h"

SIMD Calling Example

 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);
}

In NPU mode, the program prints the following information:

 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 Calling Example

 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"

// dim3 parameter when asc_vf_call is called: 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");
}

In NPU mode, the program prints the following information:

 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
......