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,表 示 无 边 界。 |
| AIV | 将Mask配 置 为 全1,表 示 所 有 数 都 参 与 计 算。 |
| AIV | 重 置loop mode的 参 数。 |
各 产 品 实 际 执 行 的 初 始 化 项 如 下:
表 1 针 对Ascend 950PR/Ascend 950DT本 接 口 执 行 的 初 始 化 项
| 状 态 类 别 | 初 始 化 内 容 |
|---|---|
| 公 共 状 态 | 原 子 累 加 状 态 |
| 公 共 状 态 | Mask工 作 模 式 |
| 公 共 状 态 | 用 于AddDeqRelu/CastDequant/Cast的s322f16场 景 的scale量 化 参 数 |
| 公 共 状 态 | CTRL寄 存 器 |
| 公 共 状 态 | 原 子 操 作 开 启 位 与 原 子 操 作 类 型 |
| AIC | Load3D接 口 调 用 时Pad填 充 的 数 值 |
| AIV | Mask |
| AIV | loop mode的 参 数 |
表 2 针 对Atlas A3 训 练 系 列 产 品/Atlas A3 推 理 系 列 产 品 本 接 口 执 行 的 初 始 化 项
| 状 态 类 别 | 初 始 化 内 容 |
|---|---|
| 公 共 状 态 | 原 子 累 加 状 态 |
| 公 共 状 态 | Mask工 作 模 式 |
| AIC | Load3D接 口 调 用 时Pad填 充 的 数 值 |
| AIC | 使 用Load3D时A1/B1的 边 界 值 |
| AIV | Mask |
表 3 针 对Atlas A2 训 练 系 列 产 品/Atlas A2 推 理 系 列 产 品 本 接 口 执 行 的 初 始 化 项
| 状 态 类 别 | 初 始 化 内 容 |
|---|---|
| 公 共 状 态 | 原 子 累 加 状 态 |
| 公 共 状 态 | Mask工 作 模 式 |
| AIC | Load3D接 口 调 用 时Pad填 充 的 数 值 |
| AIC | 使 用Load3D时A1/B1的 边 界 值 |
| AIV | Mask |
表 4 针 对Atlas 推 理 系 列 产 品 AI Core和Atlas 推 理 系 列 产 品 Vector Core本 接 口 执 行 的 初 始 化 项
| 状 态 类 别 | 初 始 化 内 容 |
|---|---|
| 公 共 状 态 | 原 子 累 加 状 态 |
函 数 原 型
__aicore__ inline void InitSocState()
参 数 说 明
无
返 回 值 说 明
无
约 束 说 明
在 实 际 运 行 中,这 些 值 可 能 被 前 序 执 行 的 算 子 修 改,若 不 调 用 该 接 口 进 行 初 始 化,非 预 期 的 值 可 能 导 致 计 算 结 果 出 现 精 度 错 误。
例 如 前 序 算 子 使 用Counter模 式 但 未 重 置 为Normal模 式,当 前 算 子 以 默 认 的Normal模 式 设 置Mask时,会 导 致Mask设 置 不 符 合 预 期,进 而 引 发 精 度 错 误。
在TPipe框 架 编 程中,初 始 化 过 程 由TPipe完 成,无 需 开 发 者 关 注;在静 态Tensor编 程的 场 景 中,用 户 必 须 在Kernel入 口 处 调 用 此 函 数 来 初 始 化AI Core状 态。
调 用 示 例
__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();
}