SK_BIND

产品支持情况

• Ascend 950PR/Ascend 950DT：不支持

• Atlas A3 训练系列产品/Atlas A3 推理系列产品：支持

• Atlas A2 训练系列产品/Atlas A2 推理系列产品：支持

• Atlas 200I/500 A2 推理产品：不支持

• Atlas 推理系列产品AI Core：不支持

• Atlas 推理系列产品Vector Core：不支持

• Atlas 训练系列产品：不支持

• Kirin X90：不支持

• Kirin 9030：不支持

功能说明

本接口为算子SuperKernel场景提供绑定原核函数和SK子函数的能力。核函数直调算子的完整适配方法请参见核函数直调算子额外适配说明。

函数原型

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

参数说明

参数名	输入/输出	描述
GF	输入	核函数函数签名。
cap	输入	uint64_t类型，用于标记当前算子的SuperKernel特性。SuperKernel框架根据这些信息调整相应的同步和融合策略。各bit位可按位或组合： • 1：early start wait flag，表示算子内使用了WaitPreTaskEnd接口。 • 2：early start set flag，表示算子内使用了SetNextTaskStart接口。 • 4：disable_dcci，表示当前算子需要SuperKernel框架关闭调用前后的DCCI ALL指令。详细说明请参见算子适配说明。 • 8：disable_batchmode_check，表示跳过当前算子的__schedmode__检查，使该算子可继续参与SuperKernel融合。 例如，同时配置early start wait flag和early start set flag时，cap取值为3。
SK0	输入	SK子函数签名。
...	输入	SK1~SK3。可提供多个SK子函数签名，包含SK0最多四个函数签名。

返回值说明

无

约束说明

• 一个核函数最多绑定四个SK子函数。
• 核函数直调算子目前仅支持在npugraph_ex后端进入SuperKernel，不支持GE图模式。

需要包含的头文件

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

#include "kernel_operator.h"

调用示例

以下示例展示SK_BIND接口的基本用法，SK子函数签名、参数结构体定义等完整规则请参见核函数直调算子额外适配说明。

#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 AddCustomArgs {
   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 AddCustomArgs *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个不同的符号
// cap取值4表示配置disable_dcci
SK_BIND(add_custom, 4, add_custom_sk<0>, add_custom_sk<1>, add_custom_sk<2>, add_custom_sk<3>);