Skip to content
版 本

InitSocState

产 品 支 持 情 况

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

功 能 说 明

头 文 件 路 径 为:"basic_api/kernel_operator_common_intf.h"

本 接 口 对AI Core的 全 局 状 态 进 行 初 始 化,包 括AIC(Cube Core)与AIV(Vector Core)的 公 共 状 态 及 各 自 特 有 状 态。由 于 不 同 产 品 的 实 现 存 在 差 异,实 际 执 行 的 初 始 化 项 也 有 所 不 同。

接 口 涉 及 的 全 部 初 始 化 操 作 汇 总 如 下:

状 态 类 别初 始 化 内 容
公 共 状 态初 始 化原 子 累 加 状 态,关 闭 数 据 搬 运 随 路 原 子 操 作 功 能。
公 共 状 态初 始 化Mask工 作 模 式为Normal模 式。
公 共 状 态初 始 化用 于AddDeqRelu/CastDequant/Cast的s322f16场 景 的scale量 化 参 数为1。
公 共 状 态初 始 化CTRL寄 存 器中 除CTRL[48]以 外 的 比 特 位 为 默 认 值。
公 共 状 态初 始 化原 子 操 作 开 启 位 与 原 子 操 作 类 型,为 无 效 的 原 子 操 作 类 型。
AIC初 始 化Load3D接 口 调 用 时Pad填 充 的 数 值为0。
AIC初 始 化使 用Load3D时A1/B1的 边 界 值为0,表 示 无 边 界。
AIVMask配 置 为 全1,表 示 所 有 数 都 参 与 计 算。
AIV重 置loop mode的 参 数

各 产 品 实 际 执 行 的 初 始 化 项 如 下:

表 1 针 对Ascend 950PR/Ascend 950DT本 接 口 执 行 的 初 始 化 项

状 态 类 别初 始 化 内 容
公 共 状 态原 子 累 加 状 态
公 共 状 态Mask工 作 模 式
公 共 状 态用 于AddDeqRelu/CastDequant/Cast的s322f16场 景 的scale量 化 参 数
公 共 状 态CTRL寄 存 器
公 共 状 态原 子 操 作 开 启 位 与 原 子 操 作 类 型
AICLoad3D接 口 调 用 时Pad填 充 的 数 值
AIVMask
AIVloop mode的 参 数

表 2 针 对Atlas A3 训 练 系 列 产 品/Atlas A3 推 理 系 列 产 品 本 接 口 执 行 的 初 始 化 项

状 态 类 别初 始 化 内 容
公 共 状 态原 子 累 加 状 态
公 共 状 态Mask工 作 模 式
AICLoad3D接 口 调 用 时Pad填 充 的 数 值
AIC使 用Load3D时A1/B1的 边 界 值
AIVMask

表 3 针 对Atlas A2 训 练 系 列 产 品/Atlas A2 推 理 系 列 产 品 本 接 口 执 行 的 初 始 化 项

状 态 类 别初 始 化 内 容
公 共 状 态原 子 累 加 状 态
公 共 状 态Mask工 作 模 式
AICLoad3D接 口 调 用 时Pad填 充 的 数 值
AIC使 用Load3D时A1/B1的 边 界 值
AIVMask

表 4 针 对Atlas 推 理 系 列 产 品 AI Core和Atlas 推 理 系 列 产 品 Vector Core本 接 口 执 行 的 初 始 化 项

状 态 类 别初 始 化 内 容
公 共 状 态原 子 累 加 状 态

函 数 原 型

C++
__aicore__ inline void InitSocState()

参 数 说 明

返 回 值 说 明

约 束 说 明

  • 在 实 际 运 行 中,这 些 值 可 能 被 前 序 执 行 的 算 子 修 改,若 不 调 用 该 接 口 进 行 初 始 化,非 预 期 的 值 可 能 导 致 计 算 结 果 出 现 精 度 错 误。

    例 如 前 序 算 子 使 用Counter模 式 但 未 重 置 为Normal模 式,当 前 算 子 以 默 认 的Normal模 式 设 置Mask时,会 导 致Mask设 置 不 符 合 预 期,进 而 引 发 精 度 错 误。

  • TPipe框 架 编 程中,初 始 化 过 程 由TPipe完 成,无 需 开 发 者 关 注;在静 态Tensor编 程的 场 景 中,用 户 必 须 在Kernel入 口 处 调 用 此 函 数 来 初 始 化AI Core状 态。

调 用 示 例

C++
__global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z)
{
    // 静 态Tensor编 程 方 式 中 需 要 开 发 者 手 动 调 用InitSocState()接 口 初 始 化 全 局 状 态 寄 存 器。
    AscendC::InitSocState();
    KernelAdd op;
    op.Init(x, y, z);
    op.Process();
}

免 责 声 明:本 站 内 容 由 asc-devkit 仓 master 分 支 自 动 编 译 生 成,属 于 持 续 开 发 版 本,可 能 存 在 缺 陷,仅 供 预 览 与 参 考。如 需 稳 定 及 商 用 资 料,请 查 阅 官 方 昇 腾 社 区