本文主要是介绍OrangePi AIpro 香橙派 昇腾 Ascend C 算子开发 与 调用 - Tiling实现 2，希望对大家解决编程问题提供一定的参考价值，需要的开发者们随着小编来一起学习吧！
OrangePi AIpro 香橙派 昇腾 Ascend C 算子开发 与 调用 - Tiling实现 2
flyfish
前置知识 1
前置知识 2
Host侧CPU和Device侧NPU的主要区别
不同的硬件资源
CPU是为了执行通用计算任务而设计的，但在处理大量的并行计算（如矩阵乘、批数据处理）时效率不高。NPU是为了加速机器学习和深度学习任务而设计的，它擅长执行大量的并行计算。NPU包含了大量的专用硬件，比如：支持矩阵计算的Cube单元，NPU中一个核可以支持一个时钟周期内完成数据量为161616、数据类型为fp16的矩阵乘法；支持向量计算的Vector单元，NPU中一个核可以支持一个时钟周期内处理128个fp16的加法。
不同的物理内存空间
Host和Device的物理内存是分离的，有时需要在Host侧内存和Device侧内存之间进行数据交换。
Ascend C的算子实现主要包含两个部分：
Host侧Tiling实现
由于NPU中AI Core内部存储无法完全容纳算子输入输出的所有数据，需要每次搬运一部分输入数据进行计算然后搬出，再搬运下一部分输入数据进行计算，这个过程就称之为Tiling。切分数据的算法称为Tiling算法或者Tiling策略。根据算子的shape等信息来确定数据切分算法相关参数（比如每次搬运的块大小，以及总共循环多少次）的计算程序，称之为Tiling实现，也叫Tiling函数（Tiling Function）。由于Tiling实现中完成的均为标量计算，AI Core并不擅长，所以我们将其独立出来放在Host侧CPU上执行。
Device侧Kernel实现
Kernel实现即算子核函数实现，在Kernel函数内部通过解析Host侧传入的Tiling结构体获取Tiling信息，根据Tiling信息控制数据搬入搬出Local Memory的流程；通过调用计算、数据搬运、内存管理、任务同步API，实现算子逻辑。其核心逻辑基本上都为计算密集型任务，需要在NPU上执行。
算子Tiling传递数据流
算子数据流
算子执行过程中涉及到Host和Device的数据交换。这里仅针对Tiling参数的传递，给出具体的数据流：Host侧Tiling算法根据算子具体输入输出的信息，完成Tiling参数的计算，存放在Tiling结构体中；将Host侧的Tiling结构体发送到Device侧，Device侧的算子获取并解析Tiling结构体，基于该信息执行后续的算子计算逻辑。
完整的代码放置最后
AddCustom算子的原型定义文件命名为add_custom.json
[{"op": "AddCustom","input_desc": [{"name": "x","param_type": "required","format": ["ND"],"type": ["fp16"]},{"name": "y","param_type": "required","format": ["ND"],"type": ["fp16"]}],"output_desc": [{"name": "z","param_type": "required","format": ["ND"],"type": ["fp16"]}]}
]
${INSTALL_DIR}/python/site-packages/bin/msopgen gen -i $HOME/sample/add_custom.json -c ai_core-<soc_version> -lan cpp -out $HOME/sample/AddCustom
${INSTALL_DIR}为CANN软件安装后文件存储路径，请根据实际环境进行替换。
-i：算子原型定义文件add_custom.json所在路径。
-c：ai_core-<soc_version>代表算子在AI Core上执行，<soc_version>为昇腾AI处理器的型号，可通过npu-smi info命令进行查询，基于同系列的AI处理器型号创建的算子工程，其基础功能能力通用。
例如soc_version设置为Ascend310P1，创建的算子工程，也可以用于开发运行于Ascend310P3上的算子。
-lan： 参数cpp代表算子基于Ascend C编程框架，使用C++编程语言开发。
实际执行的命令，截断显示多行是为了看清楚
add_custom.json所在路径是/home/HwHiAiUser/sample/add_custom.json
/usr/local/Ascend/ascend-toolkit/8.0.RC3.alpha002/python/site-packages/bin/msopgen gen
-i /home/HwHiAiUser/sample/add_custom.json
-c ai_core-Ascend310B -lan cpp
-out /home/HwHiAiUser/sample/AddCustom
执行命令后生成的文件
AddCustom
├── build.sh // 编译入口脚本
├── cmake
│ ├── config.cmake
│ ├── func.cmake
│ ├── intf.cmake
│ ├── makeself.cmake
│ └── util // 算子工程编译所需脚本及公共编译文件存放目录
├── CMakeLists.txt // 算子工程的CMakeLists.txt
├── CMakePresets.json // 编译配置项
├── framework // 算子插件实现文件目录，单算子模型文件的生成不依赖算子适配插件，无需关注
├── op_host // host侧实现文件
│ ├── add_custom_tiling.h // 算子tiling定义文件
│ ├── add_custom.cpp // 算子原型注册、shape推导、信息库、tiling实现等内容文件
│ ├── CMakeLists.txt
├── op_kernel // kernel侧实现文件
│ ├── CMakeLists.txt
│ ├── add_custom.cpp // 算子代码实现文件
└── scripts // 自定义算子工程打包相关脚本所在目录
替换为CANN软件包安装后的实际路径
修改CMakePresets.json中ASCEND_CANN_PACKAGE_PATH为CANN软件包安装路径
如果不改直接编译./build.sh就出现如下问题
在Ascend C中，Tiling策略的直接表示形式是一个C语言中的结构体(struct)，简称Tiling结构体
Tiling结构体定义在Tiling头文件(形如xxxx custom tiling.h)中，其中的每个结构体参数表示如何对输入数据进行切分，以及决定了计算过程的一些细节，结构体在host侧实例化，并通过指针传入kernel函数中
AddCustom/op_host/add_custom_tiling.h
#include "register/tilingdata_base.h"namespace optiling {
BEGIN_TILING_DATA_DEF(AddCustomTilingData)TILING_DATA_FIELD_DEF(uint32_t, size);
END_TILING_DATA_DEF;REGISTER_TILING_DATA_CLASS(AddCustom, AddCustomTilingData)
}
代码框架编写：
需要增加#ifndef…的判断条件，防止头文件的重复包含；需要包含register/tilingdata_base.h头文件，tilingdata_base.h中定义了多个用于tilingdata注册的宏。
TilingData参数设计，TilingData参数本质上是和并行数据切分相关的参数，本示例算子使用了2个tiling参数：totalLength、tileNum。totalLength是指需要计算的数据量大小，tileNum是指每个核上总计算数据分块个数。比如，totalLength这个参数传递到kernel侧后，可以通过除以参与计算的核数，得到每个核上的计算量，这样就完成了多核数据的切分。
TilingData结构定义，通过BEGIN_TILING_DATA_DEF接口定义一个TilingData的类，通过TILING_DATA_FIELD_DEF接口增加TilingData的两个字段totalLength、tileNum，通过END_TILING_DATA_DEF接口结束TilingData定义。
BEGIN_TILING_DATA_DEF(TilingData) // 注册一个tiling的类，以tiling的名字作为入参TILING_DATA_FIELD_DEF(uint32_t, totalLength); // 添加tiling字段，总计算数据量TILING_DATA_FIELD_DEF(uint32_t, tileNum); // 添加tiling字段，每个核上总计算数据分块个数
END_TILING_DATA_DEF;
注册TilingData结构，通过REGISTER_TILING_DATA_CLASS接口，注册TilingData类，和自定义算子相关联。REGISTER_TILING_DATA_CLASS第一个参数为op_type(算子类型)，本样例中传入AddCustom，第二个参数为TilingData的类名。
注册算子tilingdata类到对应的AddCustom算子REGISTER_TILING_DATA_CLASS(AddCustom, TilingData)
AddCustom/op_host/add_custom.cpp
#include "add_custom_tiling.h"
#include "register/op_def_registry.h"namespace optiling {
static ge::graphStatus TilingFunc(gert::TilingContext* context)
{AddCustomTilingData tiling;const gert::StorageShape* x1_shape = context->GetInputShape(0);int32_t data_sz = 1;for (int i = 0; i < x1_shape->GetStorageShape().GetDimNum(); i++)data_sz *= x1_shape->GetStorageShape().GetDim(i);tiling.set_size(data_sz);context->SetBlockDim(8);tiling.SaveToBuffer(context->GetRawTilingData()->GetData(), context->GetRawTilingData()->GetCapacity());context->GetRawTilingData()->SetDataSize(tiling.GetDataSize());return ge::GRAPH_SUCCESS;
}
}
AddCustom/op_kernel/add_custom.cpp
Kernel侧使用Tiling信息
Kernel侧需要接收Tiling信息时，核函数定义是这样的:
__global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling))
注意这里参数的顺序按照“输入、输出、workspace、tiling”的顺序排布，不要不要调整其顺序
#include "kernel_operator.h"
extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling) {GET_TILING_DATA(tiling_data, tiling);// TODO: user kernel impl
}
固定shape与动态shape 区别
| 文件名 | 职能 | 固定shape | 动态shape |
|---|---|---|---|
| main.cpp | 主机侧数据文件读写以及主机侧到设备侧的数据拷贝，任务下发以及同步等待等 | 读取输入参数和申请内存，调用核函数等 | 新增tiling参数的内存申请，搬运与释放逻辑 |
| add_custom.cpp | Ascend C算子核函数的实现 | shape等参数以常量展现，编译期已知 | shape等参数以入参展现，编译期未知 |
| add_custom.py | 输入数据和真值数据的生成 | 生成输入数据x和y，真值数据golden | 新增生成tiling的.bin数据文件 |
| CMakeLists.txt | 管理工程项目编译构建配置 | 无变化 | 无变化 |
| data_utils.h | 主机侧数据打印等辅助函数的实现 | 无变化 | 无变化 |
| run.sh | 集成算子运行一体化脚本 | 无变化 | 无变化 |
| add_custom_tiling.h | 定义动态shape的tiling配置 | 不涉及 | Tiling结构体与解析tiling宏函数 |
AddCustom/op_host/add_custom_tiling.h
#ifndef ADD_CUSTOM_TILING_H
#define ADD_CUSTOM_TILING_H
#include "register/tilingdata_base.h"namespace optiling {
BEGIN_TILING_DATA_DEF(TilingData)
TILING_DATA_FIELD_DEF(uint32_t, totalLength);
TILING_DATA_FIELD_DEF(uint32_t, tileNum);
END_TILING_DATA_DEF;REGISTER_TILING_DATA_CLASS(AddCustom, TilingData)
} // namespace optiling
#endif // ADD_CUSTOM_TILING_H
AddCustom/op_host/add_custom.cpp
#include "add_custom_tiling.h"
#include "register/op_def_registry.h"namespace optiling {
static ge::graphStatus TilingFunc(gert::TilingContext* context)
{AddCustomTilingData tiling;const gert::StorageShape* x1_shape = context->GetInputShape(0);int32_t data_sz = 1;for (int i = 0; i < x1_shape->GetStorageShape().GetDimNum(); i++)data_sz *= x1_shape->GetStorageShape().GetDim(i);tiling.set_size(data_sz);context->SetBlockDim(8);tiling.SaveToBuffer(context->GetRawTilingData()->GetData(), context->GetRawTilingData()->GetCapacity());context->GetRawTilingData()->SetDataSize(tiling.GetDataSize());return ge::GRAPH_SUCCESS;
}
}namespace ge {
static ge::graphStatus InferShape(gert::InferShapeContext* context)
{const gert::Shape* x1_shape = context->GetInputShape(0);gert::Shape* y_shape = context->GetOutputShape(0);*y_shape = *x1_shape;return GRAPH_SUCCESS;
}
}namespace ops {
class AddCustom : public OpDef {
public:explicit AddCustom(const char* name) : OpDef(name){this->Input("x").ParamType(REQUIRED).DataType({ge::DT_FLOAT16}).Format({ge::FORMAT_ND}).UnknownShapeFormat({ge::FORMAT_ND});this->Input("y").ParamType(REQUIRED).DataType({ge::DT_FLOAT16}).Format({ge::FORMAT_ND}).UnknownShapeFormat({ge::FORMAT_ND});this->Output("z").ParamType(REQUIRED).DataType({ge::DT_FLOAT16}).Format({ge::FORMAT_ND}).UnknownShapeFormat({ge::FORMAT_ND});this->SetInferShape(ge::InferShape);this->AICore().SetTiling(optiling::TilingFunc);this->AICore().AddConfig("ascend310b");}
};OP_ADD(AddCustom);
}
更多配置的版本
#include "add_custom_tiling.h"
#include "register/op_def_registry.h"namespace optiling {
const uint32_t BLOCK_DIM = 8;
const uint32_t TILE_NUM = 8;
static ge::graphStatus TilingFunc(gert::TilingContext *context)
{TilingData tiling;uint32_t totalLength = context->GetInputShape(0)->GetOriginShape().GetShapeSize();context->SetBlockDim(BLOCK_DIM);tiling.set_totalLength(totalLength);tiling.set_tileNum(TILE_NUM);tiling.SaveToBuffer(context->GetRawTilingData()->GetData(), context->GetRawTilingData()->GetCapacity());context->GetRawTilingData()->SetDataSize(tiling.GetDataSize());size_t *currentWorkspace = context->GetWorkspaceSizes(1);currentWorkspace[0] = 0;return ge::GRAPH_SUCCESS;
}
} // namespace optilingnamespace ge {
static graphStatus InferShape(gert::InferShapeContext *context)
{const gert::Shape *x1_shape = context->GetInputShape(0);gert::Shape *y_shape = context->GetOutputShape(0);*y_shape = *x1_shape;return GRAPH_SUCCESS;
}static graphStatus InferDataType(gert::InferDataTypeContext *context)
{const auto inputDataType = context->GetInputDataType(0);context->SetOutputDataType(0, inputDataType);return ge::GRAPH_SUCCESS;
}
} // namespace genamespace ops {
class AddCustom : public OpDef {
public:explicit AddCustom(const char *name) : OpDef(name){this->Input("x").ParamType(REQUIRED).DataType({ge::DT_FLOAT16}).Format({ge::FORMAT_ND});this->Input("y").ParamType(REQUIRED).DataType({ge::DT_FLOAT16}).Format({ge::FORMAT_ND});this->Output("z").ParamType(REQUIRED).DataType({ge::DT_FLOAT16}).Format({ge::FORMAT_ND});this->SetInferShape(ge::InferShape).SetInferDataType(ge::InferDataType);this->AICore().SetTiling(optiling::TilingFunc).AddConfig("ascend910").AddConfig("ascend310p").AddConfig("ascend310b").AddConfig("ascend910b");}
};
OP_ADD(AddCustom);
} // namespace ops
AddCustom/op_kernel/add_custom.cpp
#include "kernel_operator.h"
constexpr int32_t BUFFER_NUM = 2; // tensor num for each queueclass KernelAdd {
public:__aicore__ inline KernelAdd() {}__aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, uint32_t totalLength, uint32_t tileNum){this->blockLength = totalLength / AscendC::GetBlockNum();this->tileNum = tileNum;this->tileLength = this->blockLength / tileNum / BUFFER_NUM;xGm.SetGlobalBuffer((__gm__ DTYPE_X *)x + this->blockLength * AscendC::GetBlockIdx(), this->blockLength);yGm.SetGlobalBuffer((__gm__ DTYPE_Y *)y + this->blockLength * AscendC::GetBlockIdx(), this->blockLength);zGm.SetGlobalBuffer((__gm__ DTYPE_Z *)z + this->blockLength * AscendC::GetBlockIdx(), this->blockLength);pipe.InitBuffer(inQueueX, BUFFER_NUM, this->tileLength * sizeof(DTYPE_X));pipe.InitBuffer(inQueueY, BUFFER_NUM, this->tileLength * sizeof(DTYPE_Y));pipe.InitBuffer(outQueueZ, BUFFER_NUM, this->tileLength * sizeof(DTYPE_Z));}__aicore__ inline void Process(){int32_t loopCount = this->tileNum * BUFFER_NUM;for (int32_t i = 0; i < loopCount; i++) {CopyIn(i);Compute(i);CopyOut(i);}}private:__aicore__ inline void CopyIn(int32_t progress){AscendC::LocalTensor<DTYPE_X> xLocal = inQueueX.AllocTensor<DTYPE_X>();AscendC::LocalTensor<DTYPE_Y> yLocal = inQueueY.AllocTensor<DTYPE_Y>();AscendC::DataCopy(xLocal, xGm[progress * this->tileLength], this->tileLength);AscendC::DataCopy(yLocal, yGm[progress * this->tileLength], this->tileLength);inQueueX.EnQue(xLocal);inQueueY.EnQue(yLocal);}__aicore__ inline void Compute(int32_t progress){AscendC::LocalTensor<DTYPE_X> xLocal = inQueueX.DeQue<DTYPE_X>();AscendC::LocalTensor<DTYPE_Y> yLocal = inQueueY.DeQue<DTYPE_Y>();AscendC::LocalTensor<DTYPE_Z> zLocal = outQueueZ.AllocTensor<DTYPE_Z>();AscendC::Add(zLocal, xLocal, yLocal, this->tileLength);outQueueZ.EnQue<DTYPE_Z>(zLocal);inQueueX.FreeTensor(xLocal);inQueueY.FreeTensor(yLocal);}__aicore__ inline void CopyOut(int32_t progress){AscendC::LocalTensor<DTYPE_Z> zLocal = outQueueZ.DeQue<DTYPE_Z>();AscendC::DataCopy(zGm[progress * this->tileLength], zLocal, this->tileLength);outQueueZ.FreeTensor(zLocal);}private:AscendC::TPipe pipe;AscendC::TQue<AscendC::QuePosition::VECIN, BUFFER_NUM> inQueueX, inQueueY;AscendC::TQue<AscendC::QuePosition::VECOUT, BUFFER_NUM> outQueueZ;AscendC::GlobalTensor<DTYPE_X> xGm;AscendC::GlobalTensor<DTYPE_Y> yGm;AscendC::GlobalTensor<DTYPE_Z> zGm;uint32_t blockLength;uint32_t tileNum;uint32_t tileLength;
};extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling)
{GET_TILING_DATA(tiling_data, tiling);KernelAdd op;op.Init(x, y, z, tiling_data.totalLength, tiling_data.tileNum);op.Process();
}#ifndef ASCENDC_CPU_DEBUG
// call of kernel function
void add_custom_do(uint32_t blockDim, void *l2ctrl, void *stream, uint8_t *x, uint8_t *y, uint8_t *z,uint8_t *workspace, uint8_t *tiling)
{add_custom<<<blockDim, l2ctrl, stream>>>(x, y, z, workspace, tiling);
}
#endif
Tiling函数是在host侧实现的代码，与之相对应的，在kernel侧算子实现代码中，只需通过调用GET_TILING_DATA即可获取TilingData结构体参数，并使用具体的参数进行后续的计算。
通过Add算子举例来说明，固定shape和动态shape kernel侧算子实现的区别。
固定shape的算子例子中，TILE_NUM（每个核上总计算数据分块个数）、BLOCK_LENGTH（每个核上总计算数据大小）、TILE_LENGTH（每个分块大小）等是固定的数值。
constexpr int32_t TOTAL_LENGTH = 8 * 2048; // total length of data
constexpr int32_t USE_CORE_NUM = 8; // num of core used
constexpr int32_t BLOCK_LENGTH = TOTAL_LENGTH / USE_CORE_NUM; // length computed of each core
constexpr int32_t TILE_NUM = 8; // split data into 8 tiles for each core
constexpr int32_t BUFFER_NUM = 2; // tensor num for each queue
constexpr int32_t TILE_LENGTH = BLOCK_LENGTH / TILE_NUM / BUFFER_NUM; // each tile length is seperated to 2 part, due to double buffer
动态shape的实现中，需要在核函数中通过GET_TILING_DATA获取Tiling参数，再基于Tiling参数计算得到singleCoreSize（每个核上总计算数据大小）、tileNum（每个核上总计算数据分块个数）、singleTileLength（每个分块大小）等变量。注意，对应的算子host实现中需要定义TilingData结构体，实现并注册计算TilingData的Tiling函数。具体请参考Tiling实现。
核函数中调用GET_TILING_DATA获取TilingData的样例如下：
extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling)
{GET_TILING_DATA(tilingData, tiling);KernelAdd op;op.Init(x, y, z, tilingData.totalLength, tilingData.tileNum);if (TILING_KEY_IS(1)) {op.Process();}
}
算子类的Init函数中，使用获取到的TilingData计算得到singleCoreSize、tileNum、singleTileLength等变量的样例如下。
__aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, uint32_t totalLength, uint32_t tileNum)
{ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");this->blockLength = totalLength / GetBlockNum();this->tileNum = tileNum;ASSERT(tileNum != 0 && "tile num can not be zero!");this->tileLength = this->blockLength / tileNum / BUFFER_NUM;// ...}
这篇关于OrangePi AIpro 香橙派 昇腾 Ascend C 算子开发 与 调用 - Tiling实现 2的文章就介绍到这儿，希望我们推荐的文章对编程师们有所帮助！
