Sinh算子host侧和kernel侧代码实现分享
一、活动认证考核题目介绍
参考tensorflow的Sinh算子，实现Ascend C的Sinh算子，算子命名为SinhCustom，并完成aclnn算子调用。相关算法：sinh(x) = (exp(x) - exp(-x)) / 2.0。
本文章主要分享Ascend C下Sinh算子host侧和kernel侧代码实现分享，期望可以帮助各位初学者找到解题思路。
二、Sinh算子host侧代码实现
host侧算子实现开发包括Tiling实现、Shape推导等函数实现、算子原型注册。具体介绍如下：
- Tiling实现，计算数据切分过程相关的参数，比如每次计算的数据量大小。
- Shape推导等函数实现，根据算子的输入张量描述、算子逻辑及算子属性，推理出算子的输出张量描述，包括张量的Shape、数据类型及数据排布格式等信息。这样算子构图准备阶段就可以为所有的张量静态分配内存，避免动态内存分配带来的开销。
- 算子原型注册，除了上述函数的开发，还需要进行算子原型定义，原型定义描述了算子的输入输出、属性等信息以及算子在AI处理器上相关实现信息，算子原型注册会关联算子原型定义和上述Tiling实现、Shape推导等函数，将其组合成一个整体。
针对本题目，host侧的代码实现具体如下：
三、Sinh算子kernel侧代码实现
kernel侧的代码主要是需要通过compute函数实现Sinh算子的逻辑代码，具体如下：
#include "kernel_operator.h"
using namespace AscendC;
constexpr int32_t BUFFER_NUM = 2; // tensor num for each queue
class KernelSinh {
public:
__aicore__ inline KernelSinh() {}
__aicore__ inline void Init(GM_ADDR x, GM_ADDR y, uint32_t size)
{
ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");
this->blockLength = size / GetBlockNum();
this->tileLength = this->blockLength / BUFFER_NUM;
xGm.SetGlobalBuffer((__gm__ DTYPE_X*)x + this->blockLength * GetBlockIdx(), this->blockLength);
yGm.SetGlobalBuffer((__gm__ DTYPE_Y*)y + this->blockLength * GetBlockIdx(), this->blockLength);
pipe.InitBuffer(inQueueX, BUFFER_NUM, this->tileLength * sizeof(DTYPE_X));
pipe.InitBuffer(outQueueY, BUFFER_NUM, this->tileLength * sizeof(DTYPE_Y));
}
__aicore__ inline void Process()
{
int32_t loopCount = BUFFER_NUM;
for (int32_t i = 0; i < loopCount; i++) {
CopyIn(i);
Compute(i);
CopyOut(i);
}
}
private:
__aicore__ inline void CopyIn(int32_t progress)
{
LocalTensor<DTYPE_X> xLocal = inQueueX.AllocTensor<DTYPE_X>();
DataCopy(xLocal, xGm[progress * this->tileLength], this->tileLength);
inQueueX.EnQue(xLocal);
}
__aicore__ inline void Compute(int32_t progress)
{
LocalTensor<DTYPE_X> xLocal = inQueueX.DeQue<DTYPE_X>();
LocalTensor<DTYPE_Y> yLocal = outQueueY.AllocTensor<DTYPE_Y>();
Muls(yLocal, xLocal, (half)-1.0, this->tileLength);
Exp(yLocal, yLocal, this->tileLength);
Exp(xLocal, xLocal, this->tileLength);
Sub(yLocal, xLocal, yLocal, this->tileLength);
Muls(yLocal, yLocal, (half)0.5, this->tileLength);
outQueueY.EnQue<DTYPE_Y>(yLocal);
inQueueX.FreeTensor(xLocal);
}
__aicore__ inline void CopyOut(int32_t progress)
{
LocalTensor<DTYPE_Y> yLocal = outQueueY.DeQue<DTYPE_Y>();
DataCopy(yGm[progress * this->tileLength], yLocal, this->tileLength);
outQueueY.FreeTensor(yLocal);
}
private:
TPipe pipe;
TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX;
TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueY;
GlobalTensor<DTYPE_X> xGm;
GlobalTensor<DTYPE_Y> yGm;
uint32_t blockLength;
uint32_t tileLength;
};
extern "C" __global__ __aicore__ void sinh_custom(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) {
GET_TILING_DATA(tiling_data, tiling);
// TODO: user kernel impl
KernelSinh op;
op.Init(x, y, tiling_data.size);
op.Process();
}
#ifndef __CCE_KT_TEST__
// call of kernel function
void sinh_custom_do(uint32_t blockDim, void* l2ctrl, void* stream, uint8_t* x, uint8_t* y,
uint8_t* workspace, uint8_t* tiling)
{
sinh_custom<<<blockDim, l2ctrl, stream>>>(x, y, workspace, tiling);
}
#endif四、经验总结
本次活动学习了如何实现Ascend C算子调用，考核题目主要是需要实现Sinh算子，其中host侧和kernel侧的代码实现是其中的关键。host侧主要是需要把Tiling实现、Shape推导等函数实现、算子原型注册等完成，而kernel侧主要是需要实现compute函数，在实现该函数前，需要先理解Sinh的算子具体逻辑及实现公式，然后结合Ascend C API文档中的指令，把算子逻辑完成实现。总体来说，本次Sinh算子在逻辑实现上还是相对简单，但是收获也不少。
点赞
收藏
回复
分享
举报
Sinh算子host侧和kernel侧代码实现分享
一、活动认证考核题目介绍
参考tensorflow的Sinh算子，实现Ascend C的Sinh算子，算子命名为SinhCustom，并完成aclnn算子调用。相关算法：sinh(x) = (exp(x) - exp(-x)) / 2.0。
本文章主要分享Ascend C下Sinh算子host侧和kernel侧代码实现分享，期望可以帮助各位初学者找到解题思路。
二、Sinh算子host侧代码实现
host侧算子实现开发包括Tiling实现、Shape推导等函数实现、算子原型注册。具体介绍如下：
针对本题目，host侧的代码实现具体如下：
#include "sinh_custom_tiling.h" #include "register/op_def_registry.h" namespace optiling { static ge::graphStatus TilingFunc(gert::TilingContext* context) { SinhCustomTilingData 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 SinhCustom : public OpDef { public: explicit SinhCustom(const char* name) : OpDef(name) { this->Input("x") .ParamType(REQUIRED) .DataType({ge::DT_FLOAT16}) .Format({ge::FORMAT_ND}) .UnknownShapeFormat({ge::FORMAT_ND}); this->Output("y") .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("ascend910"); } }; OP_ADD(SinhCustom); }三、Sinh算子kernel侧代码实现
kernel侧的代码主要是需要通过compute函数实现Sinh算子的逻辑代码，具体如下：
#include "kernel_operator.h" using namespace AscendC; constexpr int32_t BUFFER_NUM = 2; // tensor num for each queue class KernelSinh { public: __aicore__ inline KernelSinh() {} __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, uint32_t size) { ASSERT(GetBlockNum() != 0 && "block dim can not be zero!"); this->blockLength = size / GetBlockNum(); this->tileLength = this->blockLength / BUFFER_NUM; xGm.SetGlobalBuffer((__gm__ DTYPE_X*)x + this->blockLength * GetBlockIdx(), this->blockLength); yGm.SetGlobalBuffer((__gm__ DTYPE_Y*)y + this->blockLength * GetBlockIdx(), this->blockLength); pipe.InitBuffer(inQueueX, BUFFER_NUM, this->tileLength * sizeof(DTYPE_X)); pipe.InitBuffer(outQueueY, BUFFER_NUM, this->tileLength * sizeof(DTYPE_Y)); } __aicore__ inline void Process() { int32_t loopCount = BUFFER_NUM; for (int32_t i = 0; i < loopCount; i++) { CopyIn(i); Compute(i); CopyOut(i); } } private: __aicore__ inline void CopyIn(int32_t progress) { LocalTensor<DTYPE_X> xLocal = inQueueX.AllocTensor<DTYPE_X>(); DataCopy(xLocal, xGm[progress * this->tileLength], this->tileLength); inQueueX.EnQue(xLocal); } __aicore__ inline void Compute(int32_t progress) { LocalTensor<DTYPE_X> xLocal = inQueueX.DeQue<DTYPE_X>(); LocalTensor<DTYPE_Y> yLocal = outQueueY.AllocTensor<DTYPE_Y>(); Muls(yLocal, xLocal, (half)-1.0, this->tileLength); Exp(yLocal, yLocal, this->tileLength); Exp(xLocal, xLocal, this->tileLength); Sub(yLocal, xLocal, yLocal, this->tileLength); Muls(yLocal, yLocal, (half)0.5, this->tileLength); outQueueY.EnQue<DTYPE_Y>(yLocal); inQueueX.FreeTensor(xLocal); } __aicore__ inline void CopyOut(int32_t progress) { LocalTensor<DTYPE_Y> yLocal = outQueueY.DeQue<DTYPE_Y>(); DataCopy(yGm[progress * this->tileLength], yLocal, this->tileLength); outQueueY.FreeTensor(yLocal); } private: TPipe pipe; TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX; TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueY; GlobalTensor<DTYPE_X> xGm; GlobalTensor<DTYPE_Y> yGm; uint32_t blockLength; uint32_t tileLength; }; extern "C" __global__ __aicore__ void sinh_custom(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) { GET_TILING_DATA(tiling_data, tiling); // TODO: user kernel impl KernelSinh op; op.Init(x, y, tiling_data.size); op.Process(); } #ifndef __CCE_KT_TEST__ // call of kernel function void sinh_custom_do(uint32_t blockDim, void* l2ctrl, void* stream, uint8_t* x, uint8_t* y, uint8_t* workspace, uint8_t* tiling) { sinh_custom<<<blockDim, l2ctrl, stream>>>(x, y, workspace, tiling); } #endif四、经验总结
本次活动学习了如何实现Ascend C算子调用，考核题目主要是需要实现Sinh算子，其中host侧和kernel侧的代码实现是其中的关键。host侧主要是需要把Tiling实现、Shape推导等函数实现、算子原型注册等完成，而kernel侧主要是需要实现compute函数，在实现该函数前，需要先理解Sinh的算子具体逻辑及实现公式，然后结合Ascend C API文档中的指令，把算子逻辑完成实现。总体来说，本次Sinh算子在逻辑实现上还是相对简单，但是收获也不少。