Ascend C算子开发指南

Ascend C的特点
C/C++原生编程:Ascend C原生支持C和C++标准规范。
屏蔽硬件差异:编程模型屏蔽了硬件差异,提高了代码的通用性。
API封装:类库API封装,既保证易用性,又兼顾高效性。
孪生调试:支持在CPU侧模拟NPU侧的行为,便于调试。
开发基本流程
环境准备:

安装CANN开发套件包,根据机器CPU架构下载对应的版本。
示例(AArch64架构):
bash
复制代码
wget -O Ascend-cann-toolkit_8.0.RC1.alpha002_linux-aarch64.run <下载链接>
chmod +x Ascend-cann-toolkit_8.0.RC1.alpha002_linux-x86_64.run
./Ascend-cann-toolkit_8.0.RC1.alpha002_linux-x86_64.run --check
sudo ./Ascend-cann-toolkit_8.0.RC1.alpha002_linux-x86_64.run --install
source /usr/local/Ascend/ascend-toolkit/set_env.sh
算子分析:

分析算子的数学表达式、输入输出数据类型和计算逻辑。
例如,Add算子的数学表达式为 $z = x + y$,输入输出数据类型为half(float16),支持的shape为(8, 2048)。
核函数开发(以Add算子为例):

获取样例代码目录quick-start,依次开发add_custom.cpp、main.cpp、gen_data.py三个文件。

核函数实现(add_custom.cpp):

cpp
复制代码
extern "C" global aicore void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z) {
KernelAdd op;
op.Init(x, y, z);
op.Process();
}

void add_custom_do(uint32_t blockDim, void* l2ctrl, void* stream, uint8_t* x, uint8_t* y, uint8_t* z) {
add_custom<<<blockDim, l2ctrl, stream>>>(x, y, z);
}
算子类实现(KernelAdd):

class KernelAdd {
public:
aicore inline KernelAdd() {}
aicore inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z) {
// 初始化代码
}
aicore inline void Process() {
// 核心处理函数
}
private:
// 各阶段函数定义
aicore inline void CopyIn(int32_t progress) {}
aicore inline void Compute(int32_t progress) {}
aicore inline void CopyOut(int32_t progress) {}
private:
TPipe pipe;
TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX, inQueueY;
TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueZ;
GlobalTensor xGm, yGm, zGm;
};
Process函数:

aicore inline void Process() {
constexpr int32_t loopCount = TILE_NUM * BUFFER_NUM;
for (int32_t i = 0; i < loopCount; i++) {
CopyIn(i);
Compute(i);
CopyOut(i);
}
}
CopyIn函数:

aicore inline void CopyIn(int32_t progress) {
LocalTensor xLocal = inQueueX.AllocTensor();
LocalTensor yLocal = inQueueY.AllocTensor();
DataCopy(xLocal, xGm[progress * TILE_LENGTH], TILE_LENGTH);
DataCopy(yLocal, yGm[progress * TILE_LENGTH], TILE_LENGTH);
inQueueX.EnQue(xLocal);
inQueueY.EnQue(yLocal);
}
Compute函数:

aicore inline void Compute(int32_t progress) {
LocalTensor xLocal = inQueueX.DeQue();
LocalTensor yLocal = inQueueY.DeQue();
LocalTensor zLocal = outQueueZ.AllocTensor();
Add(zLocal, xLocal, yLocal, TILE_LENGTH);
outQueueZ.EnQue(zLocal);
inQueueX.FreeTensor(xLocal);
inQueueY.FreeTensor(yLocal);

posted @   taixian  阅读(70)  评论(0编辑  收藏  举报
相关博文:
阅读排行:
· DeepSeek智能编程
· 精选4款基于.NET开源、功能强大的通讯调试工具
· [翻译] 为什么 Tracebit 用 C# 开发
· 腾讯ima接入deepseek-r1,借用别人脑子用用成真了~
· DeepSeek崛起:程序员“饭碗”被抢,还是职业进化新起点?
点击右上角即可分享
微信分享提示