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子函数的能力。

函数原型

// 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"头文件。

#include "kernel_operator.h"

调用示例

#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>);