SK_BIND【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit产品支持情况产品是否支持Ascend 950PR/Ascend 950DTxAtlas A3 训练系列产品 / Atlas A3 推理系列产品√Atlas A2 训练系列产品 / Atlas A2 推理系列产品√Atlas 200I/500 A2 推理产品xAtlas 推理系列产品 AI CorexAtlas 推理系列产品 Vector CorexAtlas 训练系列产品xKirin X90xKirin 9030x功能说明本接口为算子Superkernel场景提供绑定原核函数和SK子函数的能力。函数原型// GF, cap, SK0, ... #define SK_BIND(...)参数说明参数名输入/输出描述GF输入核函数函数签名。cap输入预留参数当前不使能。SuperKernel相关特性的掩码。4 表示 DCCI默认值建议使用2 表示 early start set flag1 表示 early start wait flagSK0输入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 子函数 // 模板参数仅用于实例化出不同的符号不影响函数逻辑 templateuint32_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_sk0, add_custom_sk1, add_custom_sk2, add_custom_sk3);【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考