CANN/asc-devkit：使用标准C++语法定义Tiling结构体

使用标准C语法定义Tiling结构体具体步骤在定义Tiling结构体时可以使用标准C语法定义一个POD类型Plain Old Data即与C语言兼容的数据类型。具体步骤如下。完整样例请参考标准C语法定义Tiling结构体样例。使用C语法定义Tiling结构体。[!NOTE]说明 该结构体定义所在的头文件应放置在算子工程的op_kernel目录下。由于只有该目录下的文件会被打包进算子包供在线编译场景中使用若将文件放置在其他目录中可能导致在线编译因找不到相关文件而失败。用户在使用高阶API的Tiling结构体时可通过AscendC::tiling命名空间引用kernel_tiling/kernel_tiling.h中预定义的Tiling结构体。如下代码为一个简单的Tiling结构体定义示例。#ifndef ADD_CUSTOM_TILING_H #define ADD_CUSTOM_TILING_H #include cstdint struct AddCustomTilingData { uint32_t totalLength; uint32_t tileNum; }; #endif // ADD_CUSTOM_TILING_HHost侧Tiling函数中对Tiling结构体赋值。需要包含Tiling结构体定义头文件。通过GetTilingData获取Tiling结构体指针并对其成员变量进行赋值。#include ../../op_kernel/add_custom/add_custom_tiling.h // 包含Tiling结构体定义头文件 #include register/op_def_registry.h namespace optiling { const uint32_t NUM_BLOCKS 8; const uint32_t TILE_NUM 8; static ge::graphStatus TilingFunc(gert::TilingContext *context) { // 获取Tiling结构体指针 AddCustomTilingData *tiling context-GetTilingDataAddCustomTilingData(); uint32_t totalLength context-GetInputShape(0)-GetOriginShape().GetShapeSize(); context-SetBlockDim(NUM_BLOCKS); // 对tiling的成员变量赋值 tiling-totalLength totalLength; tiling-tileNum TILE_NUM; size_t *currentWorkspace context-GetWorkspaceSizes(1); currentWorkspace[0] 0; return ge::GRAPH_SUCCESS; } } // namespace optilingKernel侧注册Tiling结构体解析Tiling数据至TilingData结构并使用。需要包含Tiling结构体定义头文件。通过REGISTER_TILING_DEFAULT或者REGISTER_TILING_FOR_TILINGKEY注册Tiling结构体通过GET_TILING_DATA解析Tiling数据至TilingData结构并使用。其中REGISTER_TILING_DEFAULT同时也用于标识使用标准C语法定义TilingData结构体。#include kernel_operator.h #include add_custom_tiling.h // 包含Tiling结构体定义头文件 ... // KernelAdd类实现 ... extern C __global__ __aicore__ void add_custom(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z, __gm__ uint8_t* workspace, __gm__ uint8_t* tiling) { REGISTER_TILING_DEFAULT(AddCustomTilingData); GET_TILING_DATA(tilingData, tiling); KernelAdd op; op.Init(x, y, z, tilingData.totalLength, tilingData.tileNum); op.Process(); }使用标准C语法定义Tiling结构体的优势相比较使用BEGIN_TILING_DATA_DEF等宏进行定义的方式该方式不仅更符合C开发者的开发习惯并且提供了强大的灵活性。支持bool类型支持数组、结构体数组及列表初始化。class A { public: bool xxx; uint32_t xxx[2][128] {0}; }; class B { public: bool xxx false; uint8_t xxx[2][2]{0}; A[10]; };不同算子可以支持定义同名但结构不同的Tiling结构体通过算子引用对应的头文件即可实现区分。这种方式允许每个算子使用符合自身需求的Tiling结构定义而不会与其他算子产生冲突。相比之下使用BEGIN_TILING_DATA_DEF等宏方式定义同名但结构不同的Tiling结构体时由于这些结构体会被注册到全局的Tiling结构体管理变量中可能导致后续通过结构体名称访问时无法准确获取当前算子实际使用的Tiling结构体从而引发未定义行为。算子Aclass TilingData { public: uint32_t length; };算子Bclass TilingData { public: uint32_t length; uint16_t coreNum; };支持自定义Tiling赋值用户可以通过访问Tiling结构体成员变量直接赋值或自定义Tiling赋值函数宏定义方式下用户仅可通过框架生成的set_xx/get_xx方法赋值/访问Tiling结构体定义class TilingData { public: uint32_t xxx1; uint32_t xxx2; uint8_t xxx3; bool xxx4; };Host侧Tiling函数#include ../op_kernel/xxx_custom_tiling.h // 包含Tiling结构体定义头文件 ... namespace optiling { static void ComputeTiling(TilingData* tiling, ...) { // 计算Tiling逻辑 ... tiling-xxx1 xxx; tiling-xxx2 xxx; tiling-xxx3 xxx; tiling-bool xxx; } static ge::graphStatus TilingFunc(gert::TilingContext *context) { ... TilingData *tiling context-GetTilingDataTilingData(); ... ComputeTiling(tiling, ...) ... return ge::GRAPH_SUCCESS; } } // namespace optiling使用约束使用标准C语法定义Tiling结构体时存在如下约束限制Tiling结构体内不支持定义成员函数因为成员函数存在Device侧和Host侧的差异Device侧的函数需要__aicore__修饰符而Tiling结构体Device侧和Host侧共用所以会在编译或执行时出现问题class TilingData { public: uint32_t xxx; __aicore__ funcA() { ... } // 错误host侧编译时不支持__aicore__修饰符会出现编译错误 void func() { ... } // 错误device侧缺少__aicore__修饰符无法执行 };Tiling结构体成员变量不支持指针、引用类型此类数据类型会导致Host侧到Device侧数据解析异常class TilingData { public: uint32_t* totalLength; // 指针场景不支持Host无法传递指针到Device uint32_t tileNum; // 引用场景不支持Host无法传递指针到Device };Tiling结构体仅支持POD类型没有虚函数、虚继承等面向对象特性也不支持模板类class A { public: uint32_t totalLength; uint32_t tileNum; }; class B: public A { public: uint32_t xxx; uint32_t xxx; }; static ge::graphStatus TilingFunc(gert::TilingContext* context) { // 错误用法 B *tiling context-GetTilingDataA(); // 不支持会触发未知问题 // 正确用法 B *tiling context-GetTilingDataB(); ...... return ge::GRAPH_SUCCESS; }GetTilingData获取的Tiling不包含初值需显式赋值或在Tiling结构体定义并调用Tiling赋值函数static ge::graphStatus TilingFunc(gert::TilingContext* context) { TilingData *tiling context-GetTilingDataTilingData(); //获取Tiling结构体此时totalLength、tileNum为0并不会带入初始值 ...... // 需显式赋值 tiling-totalLength totalLength; // 赋值Tiling结构体成员变量 tiling-tileNum TILE_NUM; // 赋值Tiling结构体成员变量 ...... return ge::GRAPH_SUCCESS; }host侧和kernel侧的Tiling结构体支持传入模板参数。由于宏函数中逗号运算符的特殊性在kernel侧宏函数REGISTER_TILING_DEFAULT或者REGISTER_TILING_FOR_TILINGKEY使用带逗号的模板类型如templateint32_t sizeA, int32_t sizeB存在编译异常因此需要使用别名方式来定义带逗号的模板类型如using size templateint32_t sizeA, int32_t sizeB。具体示例如下// 模板参数个数大于1的场景 templateint32_t sizeA, int32_t sizeB class A { public: uint32_t totalLength; uint32_t tileNum; uint32_t dataArray[sizeA]; }; // 模板参数个数等于1的场景 templateint32_t sizeA class B { public: uint32_t totalLength; uint32_t tileNum; uint32_t dataArray[sizeA]; }; // host侧可以直接传入Tiling结构体以及对应模板参数 static ge::graphStatus TilingFunc(gert::TilingContext* context) { // 模板参数个数等于1或者大于等于1的时候都可以直接传入 A3, 5 *tiling context-GetTilingDataA3,5(); B3 *tiling context-GetTilingDataB3(); ...... return ge::GRAPH_SUCCESS; } // kernel侧代码 #include kernel_operator.h #include add_custom_tiling.h // 包含Tiling结构体定义头文件 extern C __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling) { using aa A3,5; REGISTER_TILING_DEFAULT(aa); // 模板参数个数大于1时一定要用using来指定 REGISTER_TILING_FOR_TILINGKEY(TILING_KEY_VAR 2, B3); // 模板参数个数等于1时可以直接写明模板参数 ...... }如何将宏定义Tiling结构体修改为标准C语法本节介绍如何将使用BEGIN_TILING_DATA_DEF等宏进行定义的方式改造成使用标准C语法的方式。首先将之前位于op_host目录下的Tiling结构体定义头文件移至op_kernel目录下内容前后对比如下注意此时包含的头文件变化不需要再包含宏定义相关的头文件。表1两种方式对比宏定义方式标准C语法定义方式#include register/tilingdata_base.h namespace optiling { BEGIN_TILING_DATA_DEF(AddCustomTilingData) TILING_DATA_FIELD_DEF(uint32_t, totalLength); TILING_DATA_FIELD_DEF(uint32_t, tileNum); END_TILING_DATA_DEF; REGISTER_TILING_DATA_CLASS(AddCustom, AddCustomTilingData) } // namespace optiling#include cstdintstruct AddCustomTilingData { uint32_t totalLength; uint32_t tileNum; };然后修改Host侧的Tiling函数实现此时对Tiling结构体的成员变量赋值无需使用宏定义生成的set方法而是使用用户熟悉的C指针赋值方式。表2两种方式对比宏定义方式标准C语法定义方式namespace optiling { static ge::graphStatus TilingFunc(gert::TilingContext *context) { ... AddCustomTilingData tiling; tiling.set_totalLength(totalLength); // 需要使用宏定义方式生成的set方法 tiling.set_tileNum(TILE_NUM); // 需要使用宏定义方式生成的set方法 ... // 需要将局部变量保存至context上下文 tiling.SaveToBuffer(context-GetRawTilingData()-GetData(), context-GetRawTilingData()-GetCapacity()); ...return ge::GRAPH_SUCCESS;} } // namespace optiling#include ../op_kernel/add_custom/add_custom_tiling.h // 包含Tiling结构体定义头文件 ...namespace optiling { static ge::graphStatus TilingFunc(gert::TilingContext *context) { ... AddCustomTilingData *tiling context-GetTilingDataAddCustomTilingData(); ... tiling-totalLength totalLength; // 使用用户友好的C指针方式赋值成员变量 tiling-tileNum TILE_NUM; // 使用用户友好的C指针方式赋值成员变量 ...return ge::GRAPH_SUCCESS;} } // namespace optiling最后在Kernel函数入口处新增REGISTER_TILING_DEFAULT调用用于注册Tiling结构体。该注册操作的作用是告知框架用户已使用标准C语法定义Tiling结构体并明确其类型以便框架在进行Tiling数据解析时能够正确识别和使用该结构体。#include add_custom_tiling.h ... extern C __global__ __aicore__ void add_custom(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z, __gm__ uint8_t* workspace, __gm__ uint8_t* tiling) { REGISTER_TILING_DEFAULT(AddCustomTilingData); // 新增REGISTER_TILING_DEFAULT调用注册Tiling结构体 ... }创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考