SK_BIND

产 品 支 持 情 况

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

功 能 说 明

本 接 口 为 算 子Superkernel场 景 提 供 绑 定 原 核 函 数 和SK子 函 数 的 能 力。

函 数 原 型

Text

// GF, cap, SK0, ...
#define SK_BIND(...)

参 数 说 明

参 数 名	输 入/输 出	描 述
GF	输 入	核 函 数 函 数 签 名。
cap	输 入	预 留 参 数，当 前 不 开 启。 SuperKernel相 关 特 性 的 掩 码。 4 表 示 DCCI（默 认 值，建 议 使 用） 2 表 示 early start set flag 1 表 示 early start wait flag
SK0	输 入	SK子 函 数 签 名。
...	输 入	SK1~SK3 可 提 供 多 个SK子 函 数 签 名，包 含SK0最 多 四 个 函 数 签 名。

返 回 值 说 明

无

约 束 说 明

• 一 个 核 函 数 最 多 绑 定 四 个SK子 函 数。

需 要 包 含 的 头 文 件

使 用 该 接 口 需 要 包 含"kernel_operator.h"头 文 件。

Text

#include "kernel_operator.h"

调 用 示 例

Text

#include "kernel_operator.h"

// 原 普 通 kernel 保 留（用 于 非 SuperKernel 场 景）
__global__ __vector__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, uint32_t totalLength) 
{ 
    KernelAdd op; 
    op.Init(x, y, z, totalLength); 
    op.Process(); 
} 

// 规 则3：定 义 参 数 结 构 体（根 据 原 global 函 数 的 实 际 参 数 定 义）
struct GmmArgs { 
   GM_ADDR x;                      // 对 应 add_custom 的 第 一 个 参 数 
   GM_ADDR y;                      // 对 应 add_custom 的 第 二 个 参 数 
   GM_ADDR z;                      // 对 应 add_custom 的 第 三 个 参 数 
   uint32_t totalLength;           // 对 应 add_custom 的 第 四 个 参 数
};  
// 定 义 一 个 带 模 板 参 数 的 SK 子 函 数
// 模 板 参 数 仅 用 于 实 例 化 出 不 同 的 符 号，不 影 响 函 数 逻 辑
template<uint32_t splitNum>
__sk__ __vector__ void add_custom_sk(const GmmArgs *args, sk::SkSystemArgs *sysArgs/* 可 选 添 加 sysArgs 参 数*/)
{
    // 从 结 构 体 中 获 取 参 数
    GM_ADDR x = args->x;
    GM_ADDR y = args->y; 
    GM_ADDR z = args->z; 
    uint32_t totalLength = args->totalLength; 
    // 规 则5：逻 辑 与 global 函 数 一 致
    KernelAdd op;
    op.Init(x, y, z, totalLength);
    op.Process();
}
// 规 则6：使 用 SK_BIND 绑 定
// 通 过 指 定 模 板 参 数 实 例 化 出 4 个 不 同 的 符 号

SK_BIND(add_custom, 4, add_custom_sk<0>, add_custom_sk<1>, add_custom_sk<2>, add_custom_sk<3>);