ICPU_RUN_KF

产 品 支 持 情 况

产 品	是 否 支 持
Ascend 950PR/Ascend 950DT	√
Atlas A3 训 练 系 列 产 品/Atlas A3 推 理 系 列 产 品	√
Atlas A2 训 练 系 列 产 品/Atlas A2 推 理 系 列 产 品	√
Atlas 200I/500 A2 推 理 产 品	√
Atlas 推 理 系 列 产 品AI Core	√
Atlas 推 理 系 列 产 品Vector Core	x
Atlas 训 练 系 列 产 品	√
Kirin X90	√
Kirin 9030	√

功 能 说 明

头 文 件 路 径 为："tools/cpudebug/include/kern_fwk.h"。

进 行 核 函 数 的CPU侧 运 行 验 证 时，CPU调 测 总 入 口，完 成CPU侧 的 算 子 程 序 调 用。

函 数 原 型

C++

#define ICPU_RUN_KF(func, numBlocks, ...)

参 数 说 明

表 1 模 板 参 数 说 明

参 数 名 称	输 入/输 出	描 述
func	输 入	算 子 的kernel函 数 指 针。
numBlocks	输 入	算 子 的 核 心 数，corenum。
...	输 入	所 有 的 入 参 和 出 参，依 次 填 入，当 前 参 数 个 数 限 制 为32个，超 出32时 会 出 现 编 译 错 误。

返 回 值 说 明

无

约 束 说 明

• 为 了 保 留 接 口 兼 容，推 荐<<<>>>编 译 使 用。
• 除 了func、blkdim以 外，其 他 的 变 量 都 必 须 是 通 过GmAlloc分 配 的 共 享 内 存 的 指 针；传 入 的 参 数 的 数 量 和 顺 序 都 必 须 和kernel保 持 一 致。

调 用 示 例

下 面 代 码 以add_custom算 子 为 例，介 绍 算 子 核 函 数 在CPU侧 验 证 时，算 子 调 用 的 应 用 程 序 如 何 编 写。您 在 实 现 自 己 的 应 用 程 序 时，需 要 关 注 由 于 算 子 核 函 数 不 同 带 来 的 修 改，包 括 算 子 核 函 数 名，入 参 出 参 的 不 同 等，合 理 安 排 相 应 的 内 存 分 配、内 存 拷 贝 和 文 件 读 写 等，相 关API的 调 用 方 式 直 接 复 用 即 可。

1. 按 需 包 含 头 文 件，通 过ASCENDC_CPU_DEBUG宏 区 分CPU和NPU侧 需 要 包 含 的 头 文 件。

   C++

   #include "data_utils.h"
   #ifndef ASCENDC_CPU_DEBUG
   #include "acl/acl.h"
   #else
   #include "tikicpulib.h"
   extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z); // 核 函 数 声 明。
   #endif
   

2. CPU侧 运 行 验 证。完 成 算 子 核 函 数CPU侧 运 行 验 证 的 步 骤 如 下：

   图 1 CPU侧 运 行 验 证 步 骤
   

   C++

   int32_t main(int32_t argc, char* argv[])
   {
       uint32_t numBlocks = 8;
       size_t inputByteSize = 8 * 2048 * sizeof(uint16_t);
       size_t outputByteSize = 8 * 2048 * sizeof(uint16_t);
   
      // 使 用GmAlloc分 配 共 享 内 存，并 进 行 数 据 初 始 化。
       uint8_t* x = (uint8_t*)AscendC::GmAlloc(inputByteSize);
       uint8_t* y = (uint8_t*)AscendC::GmAlloc(inputByteSize);
       uint8_t* z = (uint8_t*)AscendC::GmAlloc(outputByteSize);
   
       ReadFile("./input/input_x.bin", inputByteSize, x, inputByteSize);
       ReadFile("./input/input_y.bin", inputByteSize, y, inputByteSize);
       // 矢 量 算 子 需 要 设 置 内 核 模 式 为AIV模 式。
       AscendC::SetKernelMode(KernelMode::AIV_MODE);
       // 调 用ICPU_RUN_KF调 测 宏，完 成 核 函 数CPU侧 的 调 用。
       ICPU_RUN_KF(add_custom, numBlocks, x, y, z);
       // 输 出 数 据 写 出。
       WriteFile("./output/output_z.bin", z, outputByteSize);
       // 调 用GmFree释 放 申 请 的 资 源。
       AscendC::GmFree((void *)x);
       AscendC::GmFree((void *)y);
       AscendC::GmFree((void *)z);
       return 0;
   }