Ascend C是面向算子开发场景的编程语言,它原生支持C和C++标准规范,并提供了多层接口抽象、自动并行计算等关键技术,以提高算子开发效率。
Ascend C的特点
-
C/C++原语编程
-
编程模型屏蔽硬件差异
-
类库API封装,兼顾易用与高效
-
孪生调试,可在CPU侧模拟NPU侧的行为
开发基本流程
-
环境准备
-
使用Ascend C完成Add算子核函数开发;
-
使用ICPU_RUN_KF CPU调测宏完成算子核函数CPU侧运行验证;
-
使用<<<>>>内核调用符完成算子核函数NPU侧运行验证。
下面以官方样例中的Add算子为例进行操作,开发之前请先获取样例代码目录quick-start,了解Add算子的各个目录和文件夹的具体作用,具体只需要依次开发下面的三个文件:add_custom.cpp,main.cpp,gen_data.py
Add |-- input // 存放脚本生成的输入数据目录 |-- output // 存放算子运行输出数据和真值数据的目录 |-- CMakeLists.txt // 编译工程文件 |-- add_custom.cpp // 1.算子kernel实现 |-- scripts | ├── gen_data.py // 3.输入数据和真值数据生成脚本文件 | ├── verify_result.py // 验证输出数据和真值数据是否一致的验证脚本 |-- cmake // 编译工程文件 |-- data_utils.h // 数据读入写出函数 |-- main.cpp // 2.主函数,调用算子的应用程序,含CPU域及NPU域调用 |-- run.sh // 编译运行算子的脚本
1. 环境准备
1.1 安装CANN开发套件包
-
打开CANN资源下载中心,根据机器cpu架构(uname -a来进行查看)下载对应最新的 Ascend-cann-toolkit_8.0.RC1.alpha002
# 我的机器是AArch64,因此下载安装 Ascend-cann-toolkit_8.0.RC1.alpha002_linux-aarch64.run wget -O Ascend-cann-toolkit_8.0.RC1.alpha002_linux-aarch64.run https://ascend-repo.obs.cn-east-2.myhuaweicloud.com/Milan-ASL/Milan-ASL%20V100R001C17SPC702/Ascend-cann-toolkit_8.0.RC1.alpha002_linux-aarch64.run?response-content-type=application/octet-stream
-
赋予可执行权限:
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
-
安装CANN开发套件包,并在提示时接受华为企业业务最终用户许可协议(EULA):
# root用户会默认安装在/usr/local/下,非root用户会安装到$HOME下,下面以root用户为例进行操作 ./Ascend-cann-toolkit_8.0.RC1.alpha002_linux-x86_64.run --install
-
配置CANN环境变量:
source /usr/local/Ascend/ascend-toolkit/set_env.sh
2. 算子分析
在开发Ascend C算子之前,需要对算子进行分析,明确算子的数学表达式、输入输出以及计算逻辑的实现。以Add算子为例:
-
Add算子的数学表达式为:
z = x + y
。 -
输入输出数据类型为
half
(float16),支持的shape为(8, 2048)
,format为ND。 -
核函数名称自定义为
add_custom
,参数为x
,y
,z
,分别对应输入和输出的内存地址。
3. 核函数开发(add_custom.cpp)
开发之前请先获取样例代码目录quick-start,了解Add算子的各个目录和文件夹的具体作用,具体只需要依次开发下面的三个文件:add_custom.cpp,main.cpp,gen_data.py
Add |-- input // 存放脚本生成的输入数据目录 |-- output // 存放算子运行输出数据和真值数据的目录 |-- CMakeLists.txt // 编译工程文件 |-- add_custom.cpp // 1.算子kernel实现 |-- scripts | ├── gen_data.py // 3.输入数据和真值数据生成脚本文件 | ├── verify_result.py // 验证输出数据和真值数据是否一致的验证脚本 |-- cmake // 编译工程文件 |-- data_utils.h // 数据读入写出函数 |-- main.cpp // 2.主函数,调用算子的应用程序,含CPU域及NPU域调用 |-- run.sh // 编译运行算子的脚本
3.1 核函数定义(add_custom,add_custom_do)
核函数(Kernel Function)是Ascend C算子设备侧实现的入口。在核函数中,需要为在一个核上执行的代码规定要进行的数据访问和计算操作,当核函数被调用时,多个核都执行相同的核函数代码,具有相同的参数,并行执行。
Ascend C允许用户使用核函数这种C/C++函数的语法扩展来管理设备端的运行代码,用户在核函数中进行算子类对象的创建和其成员函数的调用,由此实现该算子的所有功能。核函数是主机端和设备端连接的桥梁。
// 核函数实现 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(); // 核心处理函数,完成数据搬运与计算等核心逻辑 }
-
使用
__global__
限定符标识核函数,表示它可以在设备上执行。 -
使用
__aicore__
限定符表明核函数在AI Core上执行。 -
参数使用
GM_ADDR
宏进行修饰,表示指针指向Global Memory的内存地址。
核函数的调用需要使用特殊的内核调用符,这与普通的函数调用不同。核函数调用的语法如下:
// 核函数调用 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); }
-
blockDim
指定核函数将在多少个核上执行。 -
l2ctrl
是一个保留参数,通常设置为nullptr
。 -
stream
是aclrtStream
类型,用于维护异步操作的执行顺序。
3.2 算子类实现(KernelAdd)
class KernelAdd { public: __aicore__ inline KernelAdd(){} // 初始化函数,完成内存初始化相关操作 __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z){} // 核心处理函数,实现算子逻辑,调用私有成员函数CopyIn、Compute、CopyOut完成矢量算子的三级流水操作 __aicore__ inline void Process(){} private: // 搬入函数,完成CopyIn阶段的处理,被核心Process函数调用 __aicore__ inline void CopyIn(int32_t progress){} // 计算函数,完成Compute阶段的处理,被核心Process函数调用 __aicore__ inline void Compute(int32_t progress){} // 搬出函数,完成CopyOut阶段的处理,被核心Process函数调用 __aicore__ inline void CopyOut(int32_t progress){} private: TPipe pipe; //Pipe内存管理对象 TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX, inQueueY; //输入数据Queue队列管理对象,QuePosition为VECIN TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueZ; //输出数据Queue队列管理对象,QuePosition为VECOUT GlobalTensor<half> xGm, yGm, zGm; //管理输入输出Global Memory内存地址的对象,其中xGm, yGm为输入,zGm为输出 };
3.3 初始化函数(Init)
constexpr int32_t TOTAL_LENGTH = 8 * 2048; // 数据的总长度 constexpr int32_t USE_CORE_NUM = 8; // 使用的核心数量 constexpr int32_t BLOCK_LENGTH = TOTAL_LENGTH / USE_CORE_NUM; // 每个核心计算的长度 constexpr int32_t TILE_NUM = 8; // 每个核心将数据分割成8块 constexpr int32_t BUFFER_NUM = 2; // 每个队列的张量数量 constexpr int32_t TILE_LENGTH = BLOCK_LENGTH / TILE_NUM / BUFFER_NUM; // 由于双缓冲,将其分为两部分 __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z) { // 获取当前核心的起始索引,核心并行 xGm.SetGlobalBuffer((__gm__ half*)x + BLOCK_LENGTH * GetBlockIdx(), BLOCK_LENGTH); yGm.SetGlobalBuffer((__gm__ half*)y + BLOCK_LENGTH * GetBlockIdx(), BLOCK_LENGTH); zGm.SetGlobalBuffer((__gm__ half*)z + BLOCK_LENGTH * GetBlockIdx(), BLOCK_LENGTH); // 管道分配内存给队列,单位是字节 pipe.InitBuffer(inQueueX, BUFFER_NUM, TILE_LENGTH * sizeof(half)); pipe.InitBuffer(inQueueY, BUFFER_NUM, TILE_LENGTH * sizeof(half)); pipe.InitBuffer(outQueueZ, BUFFER_NUM, TILE_LENGTH * sizeof(half)); }
3.4 Process函数(CopyIn,Compute,CopyOut)
根据矢量编程范式,Process函数将核函数的执行分为三个基本任务:CopyIn、Compute和CopyOut。以下是这三个函数的调用方式。
__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<half> xLocal = inQueueX.AllocTensor<half>(); LocalTensor<half> yLocal = inQueueY.AllocTensor<half>(); // 将第progress个分块从全局张量复制到局部张量 DataCopy(xLocal, xGm[progress * TILE_LENGTH], TILE_LENGTH); DataCopy(yLocal, yGm[progress * TILE_LENGTH], TILE_LENGTH); // 将输入张量入队到VECIN队列 inQueueX.EnQue(xLocal); inQueueY.EnQue(yLocal); }
-
第二阶段:Compute函数的实现。
-
使用DeQue接口从VECIN队列中取出LocalTensor。
-
使用Ascend C接口Add完成矢量计算。
-
使用EnQue接口将计算结果的LocalTensor放入VECOUT类型的队列中。
-
使用FreeTensor接口释放不再使用的LocalTensor。
__aicore__ inline void Compute(int32_t progress) { // 从VECIN队列中出队输入张量 LocalTensor<half> xLocal = inQueueX.DeQue<half>(); LocalTensor<half> yLocal = inQueueY.DeQue<half>(); LocalTensor<half> zLocal = outQueueZ.AllocTensor<half>(); // 调用Add指令进行计算 Add(zLocal, xLocal, yLocal, TILE_LENGTH); // 将输出张量入队到VECOUT队列 outQueueZ.EnQue<half>(zLocal); // 释放输入张量以便重用 inQueueX.FreeTensor(xLocal); inQueueY.FreeTensor(yLocal); }
-
-
第三阶段:CopyOut函数的实现。
-
使用DeQue接口从VECOUT队列中取出LocalTensor。
-
使用DataCopy接口将LocalTensor复制到GlobalTensor上。
-
使用FreeTensor接口回收不再使用的LocalTensor。
__aicore__ inline void CopyOut(int32_t progress) { // 从VECOUT队列中出队输出张量 LocalTensor<half> zLocal = outQueueZ.DeQue<half>(); // 将第progress个分块从局部张量复制到全局张量 DataCopy(zGm[progress * TILE_LENGTH], zLocal, TILE_LENGTH); // 释放输出张量以便重用 outQueueZ.FreeTensor(zLocal); }
-
add_custom.app完整代码如下:
```c /* * 版权所有 (c) 华为技术有限公司 2022-2023。保留所有权利。 * * 函数:z = x + y * 这个示例是一个非常基础的示例,它在Ascend平台上实现了矢量加法。 * 在这个示例中: * x / y / z的长度是8*2048。 * 示例中使用的矢量核心数量是8。 * 每个核心计算的长度是2048。 * 每个核心的分块数量是8,这意味着我们在一次循环中添加2048/8=256个元素。 * * 这只是一个示范性的分块策略,实际上我们可以在一次循环中计算最多128*255个元素,对于b16类型。 */ #include "kernel_operator.h" using namespace AscendC; constexpr int32_t TOTAL_LENGTH = 8 * 2048; // 数据的总长度 constexpr int32_t USE_CORE_NUM = 8; // 使用的核心数量 constexpr int32_t BLOCK_LENGTH = TOTAL_LENGTH / USE_CORE_NUM; // 每个核心计算的长度 constexpr int32_t TILE_NUM = 8; // 每个核心将数据分成8块 constexpr int32_t BUFFER_NUM = 2; // 每个队列的张量数量 constexpr int32_t TILE_LENGTH = BLOCK_LENGTH / TILE_NUM / BUFFER_NUM; // 由于双缓冲,分为两部分 class KernelAdd { public: __aicore__ inline KernelAdd() {} __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z) { // 获取当前核心的起始索引,核心并行 xGm.SetGlobalBuffer((__gm__ half*)x + BLOCK_LENGTH * GetBlockIdx(), BLOCK_LENGTH); yGm.SetGlobalBuffer((__gm__ half*)y + BLOCK_LENGTH * GetBlockIdx(), BLOCK_LENGTH); zGm.SetGlobalBuffer((__gm__ half*)z + BLOCK_LENGTH * GetBlockIdx(), BLOCK_LENGTH); // 管道分配内存给队列,单位是字节 pipe.InitBuffer(inQueueX, BUFFER_NUM, TILE_LENGTH * sizeof(half)); pipe.InitBuffer(inQueueY, BUFFER_NUM, TILE_LENGTH * sizeof(half)); pipe.InitBuffer(outQueueZ, BUFFER_NUM, TILE_LENGTH * sizeof(half)); } __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); } } private: __aicore__ inline void CopyIn(int32_t progress) { // 从队列内存中分配张量 LocalTensor<half> xLocal = inQueueX.AllocTensor<half>(); LocalTensor<half> yLocal = inQueueY.AllocTensor<half>(); // 将第progress个分块从全局张量复制到局部张量 DataCopy(xLocal, xGm[progress * TILE_LENGTH], TILE_LENGTH); DataCopy(yLocal, yGm[progress * TILE_LENGTH], TILE_LENGTH); // 将输入张量入队到VECIN队列 inQueueX.EnQue(xLocal); inQueueY.EnQue(yLocal); } __aicore__ inline void Compute(int32_t progress) { // 从VECIN队列中出队输入张量 LocalTensor<half> xLocal = inQueueX.DeQue<half>(); LocalTensor<half> yLocal = inQueueY.DeQue<half>(); LocalTensor<half> zLocal = outQueueZ.AllocTensor<half>(); // 调用Add指令进行计算 Add(zLocal, xLocal, yLocal, TILE_LENGTH); // 将输出张量入队到VECOUT队列 outQueueZ.EnQue<half>(zLocal); // 释放输入张量以便重用 inQueueX.FreeTensor(xLocal); inQueueY.FreeTensor(yLocal); } __aicore__ inline void CopyOut(int32_t progress) { // 从VECOUT队列中出队输出张量 LocalTensor<half> zLocal = outQueueZ.DeQue<half>(); // 将第progress个分块从局部张量复制到全局张量 DataCopy(zGm[progress * TILE_LENGTH], zLocal, TILE_LENGTH); // 释放输出张量以便重用 outQueueZ.FreeTensor(zLocal); } private: TPipe pipe; // 为输入创建队列,在这种情况下深度等于缓冲区数量 TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX, inQueueY; // 为输出创建队列,在这种情况下深度等于缓冲区数量 TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueZ; GlobalTensor<half> xGm, yGm, zGm; }; // 核函数的实现 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(); } #ifndef __CCE_KT_TEST__ // 核函数的调用 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); } #endif
4. 核函数运行验证(main.cpp)
完成核函数开发后,编写host侧的调用程序进行CPU侧和NPU侧的运行验证。
4.1 主函数框架(main)
应用程序框架编写。该应用程序通过CCE_KT_TEST 宏区分代码逻辑运行于CPU侧还是NPU侧。
#include "data_utils.h" // 引入数据处理工具头文件 // 如果没有定义__CCE_KT_TEST__宏,则包含ACL头文件,否则包含TikiCPU库头文件 #ifndef __CCE_KT_TEST__ #include "acl/acl.h" // 引入AscendCL头文件,用于NPU编程 extern void add_custom_do(uint32_t coreDim, void* l2ctrl, void* stream, uint8_t* x, uint8_t* y, uint8_t* z); #else #include "tikicpulib.h" // 引入TikiCPU库头文件,用于CPU模拟NPU行为 extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z); // 声明核函数 #endif int32_t main(int32_t argc, char* argv[]) { size_t inputByteSize = 8 * 2048 * sizeof(uint16_t); // 计算输入数据的字节大小,uint16_t类型代表half精度浮点数 size_t outputByteSize = 8 * 2048 * sizeof(uint16_t); // 计算输出数据的字节大小,uint16_t类型代表half精度浮点数 uint32_t blockDim = 8; // 定义块尺寸,即在NPU上并行执行的核的数量 #ifdef __CCE_KT_TEST__ // 如果定义了__CCE_KT_TEST__宏,表示进行CPU调试,下面代码块将被编译 // 用于CPU调试的调用程序 #else // 如果没有定义__CCE_KT_TEST__宏,表示在NPU上运行,下面代码块将被编译 // NPU侧运行算子的调用程序 #endif return 0; // 程序正常退出 }
4.2 CPU侧运行验证
// 使用GmAlloc分配共享内存,并进行数据初始化 uint8_t* x = (uint8_t*)AscendC::GmAlloc(inputByteSize); uint8_t* y = (uint8_t*)AscendC::GmAlloc(inputByteSize); uint8_t* z = (uint8_t*)AscendC::GmAlloc(outputByteSize); ReadFile("./input/input_x.bin", inputByteSize, x, inputByteSize); ReadFile("./input/input_y.bin", inputByteSize, y, inputByteSize); // 矢量算子需要设置内核模式为AIV模式 AscendC::SetKernelMode(KernelMode::AIV_MODE); // 调用ICPU_RUN_KF调测宏,完成核函数CPU侧的调用 ICPU_RUN_KF(add_custom, blockDim, x, y, z); // 输出数据写出 WriteFile("./output/output_z.bin", z, outputByteSize); // 调用GmFree释放申请的资源 AscendC::GmFree((void *)x); AscendC::GmFree((void *)y); AscendC::GmFree((void *)z);
4.3 NPU侧运行验证
// AscendCL初始化 CHECK_ACL(aclInit(nullptr)); // 运行管理资源申请 aclrtContext context; int32_t deviceId = 0; CHECK_ACL(aclrtSetDevice(deviceId)); CHECK_ACL(aclrtCreateContext(&context, deviceId)); aclrtStream stream = nullptr; CHECK_ACL(aclrtCreateStream(&stream)); // 分配Host内存 uint8_t *xHost, *yHost, *zHost; uint8_t *xDevice, *yDevice, *zDevice; CHECK_ACL(aclrtMallocHost((void**)(&xHost), inputByteSize)); CHECK_ACL(aclrtMallocHost((void**)(&yHost), inputByteSize)); CHECK_ACL(aclrtMallocHost((void**)(&zHost), outputByteSize)); // 分配Device内存 CHECK_ACL(aclrtMalloc((void**)&xDevice, inputByteSize, ACL_MEM_MALLOC_HUGE_FIRST)); CHECK_ACL(aclrtMalloc((void**)&yDevice, inputByteSize, ACL_MEM_MALLOC_HUGE_FIRST)); CHECK_ACL(aclrtMalloc((void**)&zDevice, outputByteSize, ACL_MEM_MALLOC_HUGE_FIRST)); // Host内存初始化 ReadFile("./input/input_x.bin", inputByteSize, xHost, inputByteSize); ReadFile("./input/input_y.bin", inputByteSize, yHost, inputByteSize); CHECK_ACL(aclrtMemcpy(xDevice, inputByteSize, xHost, inputByteSize, ACL_MEMCPY_HOST_TO_DEVICE)); CHECK_ACL(aclrtMemcpy(yDevice, inputByteSize, yHost, inputByteSize, ACL_MEMCPY_HOST_TO_DEVICE)); // 用内核调用符<<<>>>调用核函数完成指定的运算,add_custom_do中封装了<<<>>>调用 add_custom_do(blockDim, nullptr, stream, xDevice, yDevice, zDevice); CHECK_ACL(aclrtSynchronizeStream(stream)); // 将Device上的运算结果拷贝回Host CHECK_ACL(aclrtMemcpy(zHost, outputByteSize, zDevice, outputByteSize, ACL_MEMCPY_DEVICE_TO_HOST)); WriteFile("./output/output_z.bin", zHost, outputByteSize); // 释放申请的资源 CHECK_ACL(aclrtFree(xDevice)); CHECK_ACL(aclrtFree(yDevice)); CHECK_ACL(aclrtFree(zDevice)); CHECK_ACL(aclrtFreeHost(xHost)); CHECK_ACL(aclrtFreeHost(yHost)); CHECK_ACL(aclrtFreeHost(zHost)); // AscendCL去初始化 CHECK_ACL(aclrtDestroyStream(stream)); CHECK_ACL(aclrtDestroyContext(context)); CHECK_ACL(aclrtResetDevice(deviceId)); CHECK_ACL(aclFinalize());
5. 数据生成(gen_data.py )
以固定shape的add_custom算子为例,输入数据和真值数据生成的脚本样例如下:根据算子的输入输出编写脚本,生成输入数据和真值数据。
#!/usr/bin/python3 # -*- coding:utf-8__ # 版权所有 (c) 华为技术有限公司 2022-2023。 import numpy as np # 导入numpy库,用于科学计算 # 定义生成基准数据的函数 def gen_golden_data_simple(): # 生成两个大小为[8, 2048]的随机数矩阵,范围在[-100, 100]之间,数据类型为float16 input_x = np.random.uniform(-100, 100, [8, 2048]).astype(np.float16) input_y = np.random.uniform(-100, 100, [8, 2048]).astype(np.float16) # 计算两个矩阵的和,得到基准数据(golden data),数据类型为float16 golden = (input_x + input_y).astype(np.float16) # 将生成的输入矩阵input_x和input_y分别保存到二进制文件"./input/input_x.bin"和"./input/input_y.bin" # 将基准数据golden保存到二进制文件"./output/golden.bin",用于后续结果验证 input_x.tofile("./input/input_x.bin") input_y.tofile("./input/input_y.bin") golden.tofile("./output/golden.bin") if __name__ == "__main__": # 调用函数生成基准数据 gen_golden_data_simple()
6.运行验证
完成上述文件的编写后,可以执行一键式编译运行脚本,编译和运行应用程序。
执行脚本前需要配置环境变量ASCEND_HOME_DIR,配置为CANN软件的安装路径,示例如下,请根据实际安装路径进行修改:
export ASCEND_HOME_DIR=usr/local/Ascend/ascend-toolkit/latest
脚本执行方式和脚本参数介绍如下:
-
<soc_version> :在安装昇腾AI处理器的服务器执行npu-smi info命令进行查询,在查询到的“Name”前增加Ascend信息
-
<run_mode>:表明算子以cpu模式或npu模式运行。
bash run.sh <soc_version> <run_mode>