写点什么

Ascend C 保姆级教程:我的第一份 Ascend C 代码

  • 2023-08-31
    广东
  • 本文字数:7392 字

    阅读完需:约 24 分钟

Ascend C保姆级教程:我的第一份Ascend C代码

本文分享自华为云社区《Ascend C保姆级教程:我的第一份Ascend C代码》,作者:昇腾 CANN 。


Ascend C 是昇腾 AI 异构计算架构 CANN 针对算子开发场景推出的编程语言,原生支持 C 和 C++标准规范,最大化匹配用户开发习惯;通过多层接口抽象、自动并行计算、孪生调试等关键技术,极大提高算子开发效率,助力 AI 开发者低成本完成算子开发和模型调优部署。


本文提供 Ascend C 保姆级教程,从一个简单的实例出发,带你体验 Ascend C 算子开发的基本流程。

完成实例开发之前,需要先了解一些必备的背景知识。

1. 背景知识


  • 多核并行


使用 Ascend C 开发的算子运行在 AI Core 上,AI Core 是昇腾 NPU 硬件平台的计算核心,NPU 内部有多个 AI Core。Ascend C 编程过程中会将需要处理的数据拆分同时在多个 AI Core 上运行,从而获取更高的性能。多个 AI Core 共享相同的指令代码,每个核上的运行实例唯一的区别是 block_idx 不同,开发者只需要关注单核上的处理程序,也就是核函数



  • 流水并行


上文提到,开发者只需要关注单核处理程序(核函数),那么如何实现核函数的具体逻辑呢?Ascend C 提供流水线式的编程范式,基于编程范式可以快速搭建算子实现的代码框架,实现流水并行。


流水线并行的概念和工业生产中的流水线是类似的,任务 1 完成对某片数据的处理后,将其加入到通信队列,任务 2 空闲时就会从队列中取出数据继续处理;可以类比为生产流水线中的工人只完成某一项固定工序,完成后就交由下一项工序负责人继续处理。


Ascend C 编程范式是一种流水线式的编程范式,把算子核内的处理程序,分成多个流水任务:“搬入、计算、搬出”,通过队列(Queue)完成任务间通信和同步并通过统一的内存管理模块(Pipe)管理任务间通信内存。开发者只需聚焦实现“搬入、计算、搬出”内容。



  • 孪生调试


基于 NPU 域算子的调用接口编写程序,通过毕昇编译器编译后运行,可以完成算子 NPU 域的运行验证;基于 CPU 域算子的调用接口编写程序,通过标准的 GCC 编译器进行编译后运行,并通过 GDB 通用调试工具进行单步调试,精准验证程序执行流程是否符合预期。孪生调试的能力,大大提升了算子的调试效率。下文的示例开发,仅介绍核函数 CPU 侧和 NPU 侧的运行验证,具体的调试步骤将会在后续的文章中详细介绍。

2. 开发流程


本文将引导你完成以下任务,体验 Ascend C 算子开发的基本流程。


  1. 使用 Ascend C 完成 Add 算子核函数开发

  2. 使用 ICPU_RUN_KF CPU 调测宏完成算子核函数 CPU 侧运行验证

  3. 使用<<<>>>内核调用符完成算子核函数 NPU 侧运行验证


在正式的开发之前,还需要先完成环境准备算子分析工作,开发 Ascend C 算子的基本流程如下图所示:



参考本文进行开发之前请先获取样例代码目录quick-start,该样例代码只保留了部分代码框架,核心代码在下文的指导步骤中体现。您可以在阅读本文时,将指导步骤中的代码拷贝至对应位置,即可快速完成 Ascend C 算子的开发。

3. 环境准备


  • CANN 软件安装


开发算子前,需要先准备好开发环境和运行环境,开发环境和运行环境的介绍和具体的安装步骤可参见昇腾社区文档的 CANN 软件安装指南。


  • 环境变量配置


安装 CANN 软件后,使用 CANN 运行用户编译、运行时,需要以 CANN 运行用户登录环境,执行 source ${install_path}/set_env.sh 命令设置环境变量,其中 ${install_path}为 CANN 软件的安装目录。

4. 算子分析


主要分析算子的数学表达式、输入、输出以及计算逻辑的实现,明确需要调用的 Ascend C 接口。

1. 明确算子的数学表达式及计算逻辑。


Add 算子的数学表达式为:


z = x + y


计算逻辑是:要完成 AI Core 上的数据计算,输入数据需要先从外部存储 Global Memory 搬运进 AI Core 的内部存储 Local Memory,然后使用计算接口完成两个输入参数相加,得到最终结果,再搬出到外部存储 Global Memory 上。



2. 明确输入和输出。


  • Add 算子有两个输入:x 与 y,输出为 z。

  • 本样例中算子的输入支持的数据类型为 half(float16),算子输出的数据类型与输入数据类型相同。

  • 算子输入支持 shape(8,2048),输出 shape 与输入 shape 相同。

  • 算子输入支持的 format 为:ND。


3. 确定核函数名称和参数。


  • 您可以自定义核函数名称,本样例中核函数命名为 add_custom。

  • 根据对算子输入输出的分析,确定核函数有 3 个参数 x,y,z;x,y 为输入在 Global Memory 上的内存地址,z 为输出在 Global Memory 上的内存地址。

  • 确定算子实现所需接口。实现涉及外部存储和内部存储间的数据搬运,查看 Ascend C API 参考中的数据搬移接口,需要使用 DataCopy 来实现数据搬移。本样例只涉及矢量计算的加法操作,查看 Ascend C API 参考中的矢量计算接口,初步分析可使用双目指令 Add 接口实现 x+y。计算中使用到的 Tensor 数据结构(数据操作的基础数据结构),使用 AllocTensor、FreeTensor 进行申请和释放。并行流水任务之间使用 Queue 队列完成通信和同步,会使用到 EnQue、DeQue 等接口。


通过以上分析,得到 Ascend C Add 算子的设计规格如下:


5 核函数开发


完成环境准备和初步的算子分析后,即可开始 Ascend C 核函数的开发。开发之前请先获取样例代码目录quick-start,以下核函数开发的样例代码在 add_custom.cpp 中实现。


本样例中使用多核并行计算,即把数据进行分片,分配到多个核上进行处理。Ascend C 核函数是在一个核上的处理函数,所以只处理部分数据。分配方案是:数据整体长度 TOTAL_LENGTH 为 8* 2048,平均分配到 8 个核上运行,每个核上处理的数据大小 BLOCK_LENGTH 为 2048。下文的核函数,只关注长度为 BLOCK_LENGTH 的数据应该如何处理。

5.1 核函数的定义


进行核函数的定义,并在核函数中调用算子类的 Init 和 Process 函数。请将下文代码添加至 add_custom.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();
}
复制代码


1. 使用__global__函数类型限定符来标识它是一个核函数,可以被<<<...>>>调用;使用__aicore__函数类型限定符来标识该核函数在设备端 AI Core 上执行。指针入参变量需要增加变量类型限定符__gm__,表明该指针变量指向 Global Memory 上某处内存地址为了统一表达,使用 GM_ADDR 宏来修饰入参,GM_ADDR 宏定义如下:


#define GM_ADDR __gm__ uint8_t* __restrict__
复制代码


2. 算子类的 Init 函数,完成内存初始化相关工作,Process 函数完成算子实现的核心逻辑。

5.2 算子类定义


本样例中定义 KernelAdd 算子类,其具体成员如下。请将下文代码添加至 add_custom.cpp 的“算子类实现”注释处。


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为输出
};
复制代码


内部函数的调用关系示意图如下:



由此可见除了 Init 函数完成初始化外,Process 中完成了对流水任务:“搬入、计算、搬出”的调用,开发者可以重点关注三个流水任务的实现。

5.3 Init 函数实现


初始化函数 Init 主要完成以下内容:设置输入输出 Global Tensor 的 Global Memory 内存地址,通过 Pipe 内存管理对象为输入输出 Queue 分配内存。


上文我们介绍到,本样例将数据切分成 8 块,平均分配到 8 个核上运行,每个核上处理的数据大小 BLOCK_LENGTH 为 2048。那么我们是如何实现这种切分的呢?


每个核上处理的数据地址需要在起始地址上增加 GetBlockIdx()*BLOCK_LENGTH(每个 block 处理的数据长度)的偏移来获取。这样也就实现了多核并行计算的数据切分。


以输入 x 为例,x + BLOCK_LENGTH * GetBlockIdx()即为单核处理程序中 x 在 Global Memory 上的内存偏移地址,获取偏移地址后,使用 GlobalTensor 类的 SetGlobalBuffer 接口设定该核上 Global Memory 的起始地址以及长度。具体示意图如下。



上面已经实现了多核数据的切分,那么单核上的处理数据如何进行切分?


对于单核上的处理数据,可以进行数据切块(Tiling),在本示例中,仅作为参考,将数据切分成 8 块(并不意味着 8 块就是性能最优)。切分后的每个数据块再次切分成 2 块,即可开启 double buffer,实现流水线之间的并行。


这样单核上的数据(2048 个数)被切分成 16 块,每块 TILE_LENGTH(128)个数据。Pipe 为 inQueueX 分配了两块大小为 TILE_LENGTH * sizeof(half)个字节的内存块,每个内存块能容纳 TILE_LENGTH(128)个 half 类型数据。数据切分示意图如下。


具体的初始化函数代码如下:


__aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z)
{
// 多核并行,设定当前核上Global Memory的起始地址以及长度 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为queue分配内存,单位为Bytes
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));
}
复制代码

5.4 核心处理函数实现


基于矢量编程范式,将核函数的实现分为 3 个基本任务:CopyIn,Compute,CopyOut。任务之间通过队列进行通信,交互示意图如下:



Process 函数中通过如下方式调用这三个函数。


__aicore__ inline void Process()
{
// 开启double buffer后循环次数需要乘以2
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)
{
// 1、从队列中分配Tensor
LocalTensor<half> xLocal = inQueueX.AllocTensor<half>();
LocalTensor<half> yLocal = inQueueY.AllocTensor<half>();
// 2、使用DataCopy接口将GlobalTensor数据拷贝到LocalTensor
DataCopy(xLocal, xGm[progress * TILE_LENGTH], TILE_LENGTH);
DataCopy(yLocal, yGm[progress * TILE_LENGTH], TILE_LENGTH);
// 3、将LocalTensor放入搬入数据的存放位置VecIn的Queue中
inQueueX.EnQue(xLocal);
inQueueY.EnQue(yLocal);
}
复制代码


  • Compute 函数实现。


__aicore__ inline void Compute(int32_t progress)
{
// 1、使用DeQue从VecIn中取出LocalTensor
LocalTensor<half> xLocal = inQueueX.DeQue<half>();
LocalTensor<half> yLocal = inQueueY.DeQue<half>();
LocalTensor<half> zLocal = outQueueZ.AllocTensor<half>();
// 2、调用Add指令完成双目矢量计算
Add(zLocal, xLocal, yLocal, TILE_LENGTH);
// 3、使用EnQue将计算结果LocalTensor放入到搬出数据的存放位置VECOUT的Queue中
outQueueZ.EnQue<half>(zLocal);
// 4、使用FreeTensor将释放不再使用的LocalTensor
inQueueX.FreeTensor(xLocal);
inQueueY.FreeTensor(yLocal);
}
复制代码


  • CopyOut 函数实现。


__aicore__ inline void CopyOut(int32_t progress)
{
// 1、使用DeQue接口从VecOut的Queue中取出LocalTensor
LocalTensor<half> zLocal = outQueueZ.DeQue<half>();
// 2、使用DataCopy接口将LocalTensor拷贝到GlobalTensor上
DataCopy(zGm[progress * TILE_LENGTH], zLocal, TILE_LENGTH);
// 3、使用FreeTensor将不再使用的LocalTensor进行回收
outQueueZ.FreeTensor(zLocal);
}
复制代码

6 核函数运行验证


异构计算架构中,NPU(kernel 侧)与 CPU(host 侧)是协同工作的,完成了 kernel 侧核函数开发后,即可编写 host 侧的核函数调用程序,实现从 host 侧的 APP 程序调用算子,执行计算过程。


除了上文核函数实现文件 add_custom.cpp 外,核函数的调用与验证还需要需要准备以下文件:


  • 调用算子的应用程序:main.cpp。

  • 输入数据和真值数据生成脚本文件:add_custom.py。

  • 编译 cpu 侧或 npu 侧运行的算子的编译工程文件:CMakeLists.txt。

  • 编译运行算子的脚本:run.sh。


本文仅介绍调用算子的应用程序的编写,该应用程序在 main.cpp 中体现,其他内容您可以在quick-start中直接获取。

6.1 host 侧应用程序框架编写


内置宏__CCE_KT_TEST__ 是区分运行 CPU 模式或 NPU 模式逻辑的标志,在同一个 main 函数中通过对__CCE_KT_TEST__宏定义的判断来区分 CPU 和 NPU 侧的运行程序。


int32_t main(int32_t argc, char* argv[])
{
size_t inputByteSize = 8 * 2048 * sizeof(uint16_t); // uint16_t represent half
size_t outputByteSize = 8 * 2048 * sizeof(uint16_t); // uint16_t represent half
uint32_t blockDim = 8;
#ifdef __CCE_KT_TEST__
// 用于CPU调试的调用程序


#else
// NPU侧运行算子的调用程序
#endif
return 0;
}
复制代码

6.2 CPU 运行验证


完成算子核函数 CPU 侧运行验证的步骤如下:


  1. 分配共享内存,并进行数据初始化;

  2. 调用 ICPU_RUN_KF 调测宏,完成核函数 CPU 侧的调用;

  3. 释放申请的资源。


请将下文代码添加至上面代码框架的“用于 CPU 调试的调用程序”注释处。


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);
AscendC::SetKernelMode(KernelMode::AIV_MODE);
ICPU_RUN_KF(add_custom, blockDim, x, y, z); // use this macro for cpu debug
WriteFile("./output/output_z.bin", z, outputByteSize);
AscendC::GmFree((void *)x);
AscendC::GmFree((void *)y);
AscendC::GmFree((void *)z);
复制代码

6.3 NPU 侧运行验证


完成算子核函数 NPU 侧运行验证的步骤如下:


  1. 初始化 Device 设备;

  2. 创建 Context 绑定设备;

  3. 分配 Host 内存,并进行数据初始化;

  4. 分配 Device 内存,并将数据从 Host 上拷贝到 Device 上;

  5. 用内核调用符<<<>>>调用核函数完成指定的运算;

  6. 将 Device 上的运算结果拷贝回 Host;

  7. 释放申请的资源。


请将下文代码添加至上面代码框架的“NPU 侧运行算子的调用程序”注释处。


// AscendCL初始化
CHECK_ACL(aclInit(nullptr));
// 创建Context绑定设备
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内存,并将数据从Host上拷贝到Device上
CHECK_ACL(aclrtMalloc((void**)&xDevice, inputByteSize, ACL_MEM_MALLOC_HUGE_FIRST));
CHECK_ACL(aclrtMalloc((void**)&yDevice, outputByteSize, ACL_MEM_MALLOC_HUGE_FIRST));
CHECK_ACL(aclrtMalloc((void**)&zDevice, outputByteSize, ACL_MEM_MALLOC_HUGE_FIRST));
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(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));
CHECK_ACL(aclrtDestroyStream(stream));
CHECK_ACL(aclrtDestroyContext(context));
CHECK_ACL(aclrtResetDevice(deviceId));
CHECK_ACL(aclFinalize());
复制代码

6.4 执行一键式编译运行脚本,编译和运行应用程序


脚本执行方式如下:


bash run.sh <kernel_name> <soc_version> <core_type> <run_mode>


  1. <kernel_name>表示需要运行的算子。

  2. <soc_version>表示算子运行的 AI 处理器型号。

  3. <core_type>表示在 AiCore 上或者 VectorCore 上运行。

  4. <run_mode>表示算子以 cpu 模式或 npu 模式运行。


1. CPU 模式下执行如下命令(算子运行的 AI 处理器型号以 Ascend 910 为例):


bash run.sh add_custom ascend910 AiCore cpu


运行结果如下,当前使用 md5sum 对比了所有输出 bin 文件,md5 值一致表示实际的输出数据和真值数据相符合。



2. NPU 模式下执行如下命令:bash run.sh add_custom ascend910 AiCore npu


运行结果如下,当前使用 md5sum 对比了所有输出 bin 文件,md5 值一致表示实际的输出数据和真值数据相符合。



至此,你已经完成了 Ascend C 算子开发的快速入门,更多内容请参考:《Ascend C 官方教程》


点击关注,第一时间了解华为云新鲜技术~

发布于: 2023-08-31阅读数: 2
用户头像

提供全面深入的云计算技术干货 2020-07-14 加入

生于云,长于云,让开发者成为决定性力量

评论

发布
暂无评论
Ascend C保姆级教程:我的第一份Ascend C代码_人工智能_华为云开发者联盟_InfoQ写作社区