SuperKernel融合范围标定【免费下载链接】geGEGraph Engine是面向昇腾的图编译器和执行器提供了计算图优化、多流并行、内存复用和模型下沉等技术手段加速模型执行效率减少模型内存占用。 GE 提供对 PyTorch、TensorFlow 前端的友好接入能力并同时支持 onnx、pb 等主流模型格式的解析与编译。项目地址: https://gitcode.com/cann/ge简介功能介绍SuperKernel是一种基于二进制层面的算子融合技术区别于传统的源码级融合其核心在于对已编译的二进制内核函数Kernel进行深度优化。通过智能调度与组合将多个独立的内核函数子Kernel整合为一个高效的“超级内核”SuperKernel以调用子函数的形式统一调用多个其他内核函数。相比单个算子下发SuperKernel技术显著降低了任务调度的等待时间与调度开销并能有效利用子任务间的空闲资源进一步优化整体算子的执行效率和头开销。为提升SuperKernel融合的灵活性引入了一种基于Scope标记的SuperKernel标定机制用户可手动指定需要融合到同一SuperKernel中的算子同时支持手动标记不参与融合的算子且“不融合”标记具有最高优先级。该特性目前仅支持Atlas A3 训练系列产品/Atlas A3 推理系列产品。实现原理当检测到用户进行了手动标定时系统将按照用户标定的范围进行SuperKernel融合。在SuperKernel生成策略上系统将根据其标记集合进行分组。具有完全相同标记集合的算子将被归为一组进入后续的Super Kernel融合流程。工作流程为记录标定信息系统遍历所有算子记录每个算子被用户标定的“标记集合”即哪些算子属于同一标定区域。按标记集合分组将所有算子按照其“标记集合”完全一致的原则进行分组。只有标记集合完全相同的算子才会被归入同一组。触发融合流程每组内算子进入SuperKernel的融合逻辑如下图中的sk1、sk2等。如下两个图展示了SuperKernel融合范围的标定示例标定SuperKernel融合范围图 1标记SuperKernel融合范围的示意图 使用“sk1”将op3、op4、op5标定到一个融合范围。使用“sk2”将op7、op8、op9标定到一个融合范围。op1、op2、op6由于未被纳入任何标定组不参与融合。标定SuperKernel不融合的算子图 2标记算子不融合的示意图 使用None标定的算子op3、op4、op5不参与SuperKernel的融合。使用约束标定范围内若存在不可融合的算子会生成第一段SuperKernel同时自动跳过该算子进行第二段的SuperKernel融合因此标定范围不代表最终的融合结果。融合范围标定示例通过调用aclskScopeBegin和aclskScopeEnd接口用户若传入有效scopeName可以框定算子融合进SuperKernel若传入nullptr可以框定算子不融合。该功能只在开启SuperKernel优化时有实际效果。如下代码演示了如何在NPU上利用SuperKernel技术实现一个简单的加法算子。标定算子融合到SuperKernel的关键步骤示例代码仅展示核心逻辑不可直接编译运行仅供参考#include acl/acl.h #include super_kernel/super_kernel.h ... // 原始kernel函数add_custom __global__ __vector__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, AddCustomTilingData tiling) { KernelAdd op; op.Init(x, y, z, tiling.totalLength, tiling.tileNum); op.Process(); } // 定义参数结构体用于封装传递给SuperKernel的参数 struct ArgsStruct { GM_ADDR x; GM_ADDR y; GM_ADDR z; AddCustomTilingData tiling; }; // 定义一个带模板参数的SuperKernel子函数add_custom_sk templateuint32_t splitNum __sk__ __vector__ void add_custom_sk(const ArgsStruct *args, sk::SkSystemArgs *sysArgs) { // 从结构体获取参数 GM_ADDR x args-x; GM_ADDR y args-y; GM_ADDR z args-z; AddCustomTilingData tiling args-tiling; uint16_t blockNum sysArgs-skNumBlocks; uint16_t blockIdx sysArgs-skBlockIdx; // 逻辑与原kernel的global函数一致 KernelAdd op; op.Init(x, y, z, tiling.totalLength, tiling.tileNum); op.Process(); } // 使用SK_BIND绑定将add_custom kernel与多个SK子函数绑定通过指定模板参数实例化出4个不同的符号 SK_BIND(add_custom, 4, add_custom_sk0, add_custom_sk1, add_custom_sk2, add_custom_sk3); // 主函数 int main() { // 算子数据准备 constexpr uint32_t totalLength 8 * 2048; constexpr float valueX 1.2f; constexpr float valueY 2.3f; std::vectorfloat x(totalLength, valueX); std::vectorfloat y(totalLength, valueY); constexpr uint32_t numBlocks 8; uint32_t totalLength x.size(); size_t totalByteSize totalLength * sizeof(float); int32_t deviceId 0; aclrtStream stream nullptr; AddCustomTilingData tiling {totalLength, 8}; uint8_t *xHost reinterpret_castuint8_t *(x.data()); // printf(0-host: x0:%d \n, xHost[0]); uint8_t *yHost reinterpret_castuint8_t *(y.data()); uint8_t *zHost nullptr; uint8_t *xDevice nullptr; uint8_t *yDevice nullptr; uint8_t *zDevice nullptr; // 内存分配 aclrtMallocHost((void **)(zHost), totalByteSize); aclrtMalloc((void **)xDevice, totalByteSize, ACL_MEM_MALLOC_HUGE_FIRST); aclrtMalloc((void **)yDevice, totalByteSize, ACL_MEM_MALLOC_HUGE_FIRST); aclrtMalloc((void **)zDevice, totalByteSize, ACL_MEM_MALLOC_HUGE_FIRST); // 数据传输 aclrtMemcpy(xDevice, totalByteSize, xHost, totalByteSize, ACL_MEMCPY_HOST_TO_DEVICE); aclrtMemcpy(yDevice, totalByteSize, yHost, totalByteSize, ACL_MEMCPY_HOST_TO_DEVICE); // 初始化 aclInit(nullptr); aclrtSetDevice(deviceId); aclrtCreateStream(stream); // 构建图 aclmdlRICaptureBegin(stream, ACL_MODEL_RI_CAPTURE_MODE_GLOBAL); aclmdlRI modelRI; // 标定SuperKernel融合起始位置 aclskScopeBegin(sk1, stream); add_customnumBlocks, nullptr, stream(xDevice, yDevice, zDevice, tiling); // 标定SuperKernel融合结束位置 aclskScopeEnd(sk1, stream); aclmdlRICaptureEnd(stream, modelRI); // 开启SuperKernel优化 aclskOptimize(modelRI, nullptr); // 执行图 aclmdlRIExecuteAsync(modelRI, stream); // 获取图执行结果 aclrtSynchronizeStream(stream); aclrtMemcpy(zHost, totalByteSize, zDevice, totalByteSize, ACL_MEMCPY_DEVICE_TO_HOST); std::vectorfloat z((float *)zHost, (float *)(zHost totalLength)); // 资源释放 aclrtFree(xDevice); aclrtFree(yDevice); aclrtFree(zDevice); aclrtFreeHost(zHost); aclrtDestroyStream(stream); aclrtResetDevice(deviceId); aclFinalize(); }上述acl接口详细说明请参见《Runtime运行时 API》。标定算子不融合到SuperKernel的关键步骤示例以下代码仅展示核心逻辑不可直接编译运行仅供参考#include acl/acl.h #include super_kernel/super_kernel.h ... int main() { // 算子数据准备 ... // 初始化 aclInit(nullptr); aclrtSetDevice(deviceId); aclrtCreateStream(stream); // 构建图 aclmdlRICaptureBegin(stream, ACL_MODEL_RI_CAPTURE_MODE_GLOBAL); aclmdlRI modelRI; // 标定SuperKernel融合起始位置传入nullptr表示不融合 aclskScopeBegin(nullptr, stream); add_customnumBlocks, nullptr, stream(xDevice, yDevice, zDevice, tiling); // 标定SuperKernel融合结束位置传入nullptr表示不融合 aclskScopeEnd(nullptr, stream); aclmdlRICaptureEnd(stream, modelRI); // 开启SuperKernel优化 aclskOptimize(modelRI, nullptr); // 执行图 aclmdlRIExecuteAsync(modelRI, stream); // 获取图执行结果 ... // 资源释放 aclrtFree(xDevice); aclrtFree(yDevice); aclrtFree(zDevice); aclrtFreeHost(zHost); aclrtDestroyStream(stream); aclrtResetDevice(deviceId); aclFinalize(); }上述acl接口详细说明请参见《Runtime运行时 API》。【免费下载链接】geGEGraph Engine是面向昇腾的图编译器和执行器提供了计算图优化、多流并行、内存复用和模型下沉等技术手段加速模型执行效率减少模型内存占用。 GE 提供对 PyTorch、TensorFlow 前端的友好接入能力并同时支持 onnx、pb 等主流模型格式的解析与编译。项目地址: https://gitcode.com/cann/ge创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考