Async

产 品 支 持 情 况

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

功 能 说 明

头 文 件 路 径 为："basic_api/kernel_operator_utils_intf.h"。

基 于分 离 模 式（AIC和AIV分 离）开 发 融 合 算 子 时，算 子 逻 辑 中 通 常 同 时 包 含AIV和AIC的 处 理 逻 辑，此 时 需 要 使 用Ascend C提 供 的 宏ASCEND_IS_AIV/ASCEND_IS_AIC实 现 如 下 硬 件 条 件 分 支 来 对AIV和AIC的 代 码 进 行 隔 离：

C++

if ASCEND_IS_AIV {
// AIV处 理 逻 辑
}
if ASCEND_IS_AIC {
// AIC处 理 逻 辑
}

Async通 过 模 板 函 数 的 方 式 对 这 种 隔 离 模 式 进 行 了 封 装，提 供 了 一 个 统 一 的 接 口，用 于 在 不 同 执 行 单 元（AIC或AIV）下 执 行 特 定 函 数，从 而 避 免 在 代 码 中 使 用 硬 件 条 件 分 支。

函 数 原 型

C++

template <EngineType engine, auto funPtr, class... Args>
__aicore__ void Async(Args... args)

参 数 说 明

表 1 模 板 参 数 说 明

参 数 名	描 述
engine	EngineType枚 举 类 型，可 取 值 如 下，对 应 不 同 的 硬 件 执 行 单 元： • AIC • AIV
funPtr	函 数 指 针，指 定 要 执 行 的 函 数，函 数 签 名 和 参 数 类 型 由class... Args决 定。
class... Args	可 变 参 数 模 板，表 示 函 数 参 数 的 类 型 列 表，用 于 传 递 给funPtr。

表 2 参 数 说 明

参 数 名	输 入/输 出	描 述
Args... args	输 入	与class... Args对 应 的 参 数 列 表，表 示 传 递 给funPtr的 实 际 参 数。

返 回 值 说 明

无

约 束 说 明

无

调 用 示 例

C++

__aicore__ inline void cubeProcess(KernelMmad &op, GM_ADDR A, GM_ADDR B, GM_ADDR c)
{
    op.InitAIC(A, B, c);
    op.ProcessAIC();
}
__aicore__ inline void vectorProcess(KernelMmad &op, GM_ADDR a, GM_ADDR b, GM_ADDR A, GM_ADDR B, GM_ADDR c)
{
    op.InitAIV(a, b, A, B, c);
    op.ProcessAIV();
}
__global__ __mix__(1,2) void mmad_custom(GM_ADDR a, GM_ADDR b, GM_ADDR A, GM_ADDR B, GM_ADDR c)
{
    AscendC::InitSocState();
    KernelMmad op;
    AscendC::Async<AscendC::EngineType::AIC, cubeProcess>(op, A, B, c);
    AscendC::Async<AscendC::EngineType::AIV, vectorProcess>(op, a, b, A, B, c);
}