Ascend C
搜 索

Appearance

版 本
本 页 内 容

printf

产 品 支 持 情 况​

  • 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接 口 打 印 相 关 内 容。

注 意

printf（PRINTF）接 口 打 印 功 能 会 对 算 子 实 际 运 行 的 性 能 带 来 一 定 影 响，通 常 在 调 测 阶 段 使 用。开 发 者 可 以 按 需 通 过 设 置ASCENDC_DUMP=0的 方 式 关 闭 打 印 功 能。

函 数 原 型​

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

以 下 接 口 为simd_vf中 所 使 用 的printf接 口，仅 支 持Ascend 950PR/Ascend 950DT。

Text
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 VF中printf功 能 需 要 占 用 额 外 的Global Memory空 间 用 于 数 据 缓 存，缓 存 空 间 大 小 默 认 为2MB。您 可 以 通 过acl.json中 的"simt_printf_fifo_size"字 段 进 行 配 置，配 置 范 围 最 小 为1MB，最 大 为64MB。当 打 印 数 据 量 较 大 时，建 议 增 加 缓 存 空 间。

  • printf在SIMT VF中 调 用，会 占 用SIMT栈 空 间，请 合 理 控 制 调 用printf次 数，防 止SIMT栈 溢 出。

  • 在 仿 真 环 境 下，使 用printf接 口 会 增 加 算 子 运 行 时 间，通 过 在VF代 码 中 判 断 线 程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二 进 制 文 件 解 析 为 用 户 可 读 内 容，命 令 格 式 如 下。

      Text
      show_kernel_debug_data  bin_file  output_dir

  • 由 于 在simd vf中 不 能 直 接 访 问gm地 址，因 此 使 用simd vf中 的printf时，需 要 先 定 义 一 个ubuf的 字 符 串，该 字 符 串 即 为 原 来 的fmt变 量。

    Text
    __ubuf__ const char* fmt = "simd vf: int=%d, uint=%u, float=%f, string=%s\n";

需 要 包 含 的 头 文 件​

使 用 该 接 口 需 要 包 含"utils/debug/asc_printf.h"头 文 件。

Text
#include "utils/debug/asc_printf.h"

SIMD调 用 示 例​

Text
#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模 式 下，程 序 运 行 时 打 印 效 果 如 下：

Text
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调 用 示 例​

Text
#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模 式 下，程 序 运 行 时 打 印 效 果 如 下：

Text
simd vf: int=1, uint=2, float=5.000000, string=AscendC
simd vf: int=1, uint=2, float=5.000000, string=AscendC
simd vf: int=1, uint=2, float=5.000000, string=AscendC
......

SIMT示 例​

Text
#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模 式 下，程 序 运 行 时 打 印 效 果 如 下：

Text
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 仓 master 分 支 自 动 编 译 生 成，属 于 持 续 开 发 版 本，可 能 存 在 缺 陷，仅 供 预 览 与 参 考。如 需 稳 定 及 商 用 资 料，请 查 阅 官 方 昇 腾 社 区。

Copyright Huawei Technologies Co., Ltd.