设 置Kernel类 型
产 品 支 持 情 况
功 能 说 明
用 于 用 户 自 定 义 设 置kernel类 型，控 制 算 子 执 行 时 只 启 动 该 类 型 的 核，避 免 启 动 不 需 要 工 作 的 核，缩 短 核 启 动 开 销。
函 数 原 型
设 置 全 局 默 认 的kernel type，对 所 有 的tiling key生 效。
当 前 支 持 在 自 定 义 算 子 工 程 和Kernel直 调 工 程 中 使 用。
TextKERNEL_TASK_TYPE_DEFAULT(value)设 置 某 一 个 具 体 的tiling key对 应 的kernel type。
当 前 仅 支 持 在 自 定 义 算 子 工 程 中 使 用。
TextKERNEL_TASK_TYPE(key, value)
参 数 说 明
表 1 参 数 说 明
设 置 的kernel类 型，可 选 值 范 围，kernel类 型 具 体 说 明 请 参 考表2。不 同 硬 件 架 构 支 持 的 参 数 取 值 不 同，具 体 支 持 的 参 数 取 值 请 参 考kernel type取 值 约 束。 enum KernelMetaType {
KERNEL_TYPE_AIV_ONLY,
KERNEL_TYPE_AIC_ONLY,
KERNEL_TYPE_MIX_AIV_1_0,
KERNEL_TYPE_MIX_AIC_1_0,
KERNEL_TYPE_MIX_AIC_1_1,
KERNEL_TYPE_MIX_AIC_1_2,
KERNEL_TYPE_AICORE,
KERNEL_TYPE_VECTORCORE,
KERNEL_TYPE_MIX_AICORE,
KERNEL_TYPE_MIX_VECTOR_CORE,
KERNEL_TYPE_MAX
}; |
表 2 kernel type取 值 说 明
算 子 执 行 时 仅 启 动AI Core上 的Vector核：比 如 用 户 在host侧 设 置numBlocks为10，则 会 启 动10个Vector核。 | |
算 子 执 行 时 仅 启 动AI Core上 的Cube核：比 如 用 户 在host侧 设 置numBlocks为10，则 会 启 动10个Cube核。 | |
AIC、AIV混 合 场 景 下，使 用 了多 核 控 制 相 关 指 令时，设 置 核 函 数 的 类 型 为MIX AIV:AIC 1:0（带 有 硬 同 步），算 子 执 行 时 仅 会 启 动AI Core上 的Vector核，比 如 用 户 在host侧 设 置numBlocks为10，则 会 启 动10个Vector核。 硬 同 步 的 概 念 解 释 如 下：当 不 同 核 之 间 操 作 同 一 块 全 局 内 存 且 可 能 存 在 读 后 写、写 后 读 以 及 写 后 写 等 数 据 依 赖 问 题 时，通 过 调 用SyncAll()函 数 来 插 入 同 步 语 句 来 避 免 上 述 数 据 依 赖 时 可 能 出 现 的 数 据 读 写 错 误 问 题。目 前 多 核 同 步 分 为 硬 同 步 和 软 同 步，硬 同 步 是 利 用 硬 件 自 带 的 全 核 同 步 指 令 由 硬 件 保 证 多 核 同 步。 | |
AIC、AIV混 合 场 景 下，使 用 了多 核 控 制 相 关 指 令时，设 置 核 函 数 的 类 型 为MIX AIC:AIV 1:0（带 有 硬 同 步），算 子 执 行 时 仅 会 启 动AI Core上 的Cube核，比 如 用 户 在host侧 设 置numBlocks为10，则 会 启 动10个Cube核。 | |
AIC、AIV混 合 场 景 下，设 置 核 函 数 的 类 型 为MIX AIC:AIV 1:1，算 子 执 行 时 会 同 时 启 动AI Core上 的Cube核 和Vector核，比 如 用 户 在host侧 设 置numBlocks为10，则 会 启 动10个Cube核 和10个Vector核。 | |
AIC、AIV混 合 场 景 下，设 置 核 函 数 的 类 型 为MIX AIC:AIV 1:2，算 子 执 行 时 会 同 时 启 动AI Core上 的Cube核 和Vector核，比 如 用 户 在host侧 设 置numBlocks为10，则 会 启 动10个Cube核 和20个Vector核。 | |
算 子 执 行 时 仅 会 启 动AI Core，比 如 用 户 在host侧 设 置numBlocks为5，则 会 启 动5个AI Core。 | |
基 于Ascend C开 发 的 矢 量 计 算 相 关 的 算 子 可 以 运 行 在Vector Core上，调 用 本 接 口 传 入 该 参 数 用 于 启 用Vector Core。 启 用Vector Core后，算 子 执 行 时 会 同 时 启 动AI Core和Vector Core，用 于 并 行 计 算。比 如 用 户 在host侧 设 置numBlocks为10，则 会 启 动 总 数 为10的AI Core和Vector Core。 需 要 注 意 的 是，通 过SetBlockDim设 置 核 数 时，需 要 大 于AI Core的 核 数，否 则 不 会 启 动VectorCore。 |
约 束 说 明
- Ascend 950PR/Ascend 950DT，支 持KERNEL_TYPE_AIV_ONLY、 KERNEL_TYPE_AIC_ONLY、KERNEL_TYPE_MIX_AIV_1_0、KERNEL_TYPE_MIX_AIC_1_0、KERNEL_TYPE_MIX_AIC_1_1、KERNEL_TYPE_MIX_AIC_1_2。
- Atlas A3 训 练 系 列 产 品/Atlas A3 推 理 系 列 产 品，支 持KERNEL_TYPE_AIV_ONLY、 KERNEL_TYPE_AIC_ONLY、KERNEL_TYPE_MIX_AIV_1_0、KERNEL_TYPE_MIX_AIC_1_0、KERNEL_TYPE_MIX_AIC_1_1、KERNEL_TYPE_MIX_AIC_1_2。
- Atlas A2 训 练 系 列 产 品/Atlas A2 推 理 系 列 产 品，支 持KERNEL_TYPE_AIV_ONLY、 KERNEL_TYPE_AIC_ONLY、KERNEL_TYPE_MIX_AIV_1_0、KERNEL_TYPE_MIX_AIC_1_0、KERNEL_TYPE_MIX_AIC_1_1、KERNEL_TYPE_MIX_AIC_1_2。
- Atlas 推 理 系 列 产 品，支 持KERNEL_TYPE_AICORE、KERNEL_TYPE_MIX_VECTOR_CORE。
KERNEL_TASK_TYPE优 先 级 高 于KERNEL_TASK_TYPE_DEFAULT，同 时 设 置 了 全 局kernel type和 某 一 个tiling key的kernel type，该tiling key的kernel type以KERNEL_TASK_TYPE设 置 的 为 准。
没 有 设 置 全 局 默 认kernel type的 情 况 下，如 果 开 发 者 只 为 其 中 的 某 几 个tiling key设 置kernel type，即 部 分tiling key没 有 设 置kernel type，会 导 致 算 子kernel编 译 报 错。
当 设 置 具 体 的kernel task type时，用 户 的 算 子 实 现 需 要 与kernel type相 匹 配。比 如 用 户 设 置kernel type为KERNEL_TYPE_MIX_AIC_1_2，则 算 子 内 部 实 现 应 与 核 配 比AIC:AIV为1:2相 对 应；若 用 户 设 置kernel type为KERNEL_TYPE_AIC_ONLY， 则 算 子 内 部 实 现 应 该 为 纯cube逻 辑，不 应 该 存 在vector部 分 的 逻 辑。其 他 的kernel type类 似。
当 纯cube或 者 纯vec算 子 强 制 设 定kernel type为MIX类 型 时，workspace的 大 小 不 能 设 置 为0，需 要 设 置 一 个 大 于0的 值（比 如16、32等）。
使 用Tiling模 板 编 程时，需 要 通 过ASCENDC_TPL_KERNEL_TYPE_SEL设 置Kernel类 型 即 可，无 需 再 通 过 该 接 口 进 行 设 置，本 接 口 不 生 效。
调 用 示 例
示 例 一：启 用VectorCore样 例
完 成 算 子kernel侧 开 发 时，需 要 通 过 本 接 口 启 用Vector Core，算 子 执 行 时 会 同 时 启 动AI Core和Vector Core， 此 时AI Core会 当 成Vector Core使 用。示 例 如 下：
Textextern "C" __global__ __aicore__ void add_custom(__gm__ uint8_t *x, __gm__ uint8_t *y, __gm__ uint8_t *z, __gm__ uint8_t *workspace, __gm__ uint8_t *tiling) { GET_TILING_DATA(tilingData, tiling); if (workspace == nullptr) { return; } KernelAdd op; op.Init(x, y, z, tilingData.numBlocks, tilingData.totalLength, tilingData.tileNum); KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_VECTOR_CORE); // 启 用VectorCore if (TILING_KEY_IS(1)) { op.Process1(); } else if (TILING_KEY_IS(2)) { op.Process2(); } // ... }完 成 算 子host侧Tiling开 发 时，设 置 的numBlocks代 表 的 是AI Core和Vector Core的 总 数，比 如 用 户 在host侧 设 置numBlocks为10，则 会 启 动 总 数 为10的AI Core和Vector Core；为 保 证 启 动Vector Core，设 置 数 值 应 大 于AI Core的 核 数。您 可 以 通 过GetCoreNumAic接 口 获 取AI Core的 核 数，GetCoreNumVector接 口 获 取Vector Core的 核 数。 如 下 代 码 片 段，展 示 了numBlocks的 设 置 方 法，此 处 设 置 为AI Core和Vector Core的 总 和，表 示 所 有AI Core和Vector Core都 启 动。
Text// 配 套 的host侧tiling函 数 示 例： ge::graphStatus TilingFunc(gert::TilingContext* context) { // 启 用VectorCore，将numBlocks置 为AI Core中vector核 数 + Vector Core中 的vector核 数 auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfo); auto totalCoreNum = ascendcPlatform.GetCoreNumAiv(); // ASCENDXXX请 替 换 为 实 际 的 版 本 型 号 if (ascendcPlatform.GetSocVersion() == platform_ascendc::SocVersion::ASCENDXXX) { totalCoreNum = totalCoreNum + ascendcPlatform.GetCoreNumVector(); } context->SetBlockDim(totalCoreNum); }
示 例 二：设 置 某 一 个 具 体 的tiling key对 应 的kernel type。如 下 代 码 为 伪 代 码 ，不 可 直 接 运 行。
Textextern "C" __global__ __aicore__ void add_custom(__gm__ uint8_t *x, __gm__ uint8_t *y, __gm__ uint8_t *z, __gm__ uint8_t *workspace, __gm__ uint8_t *tiling) { GET_TILING_DATA(tilingData, tiling); if (workspace == nullptr) { return; } KernelAdd op; op.Init(x, y, z, tilingData.numBlocks, tilingData.totalLength, tilingData.tileNum); KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY); // 设 置 默 认 的kernel类 型 为 纯AIV类 型 if (TILING_KEY_IS(1)) { KERNEL_TASK_TYPE(1, KERNEL_TYPE_MIX_AIV_1_0); // 设 置tiling key=1对 应 的kernel类 型 为MIX AIV 1:0 op.Process1(); } else if (TILING_KEY_IS(2)) { KERNEL_TASK_TYPE(2, KERNEL_TYPE_AIV_ONLY); // 设 置tiling key=2对 应 的kernel类 型 为 纯AIV类 型 op.Process2(); } // ... } // 配 套 的host侧tiling函 数 示 例： ge::graphStatus TilingFunc(gert::TilingContext* context) { // ... if (context->GetInputShape(0) > 10) { context->SetTilingKey(1); } else if (some condition) { context->SetTilingKey(2); } }