本文分享自华为云社区《昇腾Ascend C编程入门教程》，作者：昇腾 CANN 。
2023 年 5 月 6 日，在昇腾 AI 开发者峰会上，华为正式发布了面向算子开发场景的昇腾 Ascend C 编程语言。Ascend C 原生支持 C/C++编程规范，通过多层接口抽象、并行编程范式、孪生调试等技术，极大提高了算子的开发效率，帮助 AI 开发者低成本完成算子开发和模型调优部署。
昇腾 AI 软硬件基础
和 CUDA 开发的算子运行在 GPU 上一样，基于 Ascend C 开发的算子，可以通过异构计算架构 CANN（Compute Architecture for Neural Networks）运行在昇腾 AI 处理器（可简称 NPU）上。CANN 是使能昇腾 AI 处理器的一个软件栈，通过软硬件协同优化，能够充分发挥昇腾 AI 处理器的强大算力。从下面的架构图可以清楚的看到，使用 Ascend C 编程语言开发的算子通过编译器编译和运行时调度，最终运行在昇腾 AI 处理器上。
我们知道，通用计算就是我们常写的一些在 CPU 上运行的计算，它擅长逻辑控制和串行计算，而 AI 计算相对通用计算来说，更擅长并行计算，可支持大规模的计算密集型任务。如下面左图所示，做一个矩阵乘，使用 CPU 计算需要三层 for 循环，而右图在昇腾 AI 处理器上使用 vector 计算单元，只需要两层 for 循环，最小计算代码能同时计算多个数据的乘加，更近一步，如果使用 Cube 计算单元，只需要一条语句就能完成一个矩阵乘的计算，这就是我们所说的 SIMD（单指令多数据）。因此，我们通常使用 AI 处理器来进行大量的并行计算。
NPU 不能独立运行，需要与 CPU 协同工作，可以看成是 CPU 的协处理器，CPU 负责整个操作系统运行，管理各类资源并进行复杂的逻辑控制，而 NPU 主要负责并行计算任务。在基于 CPU+NPU 的异构计算架构中，NPU 与 CPU 通过 PCIe 总线连接在一起来协同工作，CPU 所在位置称为主机端（host），而 NPU 所在位置称为设备端（device），示意图如下：
这里再详细介绍一下昇腾 AI 处理器。昇腾 AI 处理器有不同的型号和产品形态，小到模块、加速卡，大到服务器、集群。昇腾 AI 处理器里面最核心的部件是 AI Core，有多个，是神经网络加速的计算核心，每一个 AI Core 就相当于我们大家平时理解的多核 cpu 里的每个核，使用 Ascend C 编程语言开发的算子就运行在 AI Core 上，因为核心的神经网络计算的加速都来源于 AI Core 的算力。
AI Core 内部的并行计算架构抽象如下图所示：
这个并行计算架构抽象核心包含了几个大的部件，AI Core 外面有一个 Gobal Memory，是多个 AI Core 共享的，在 AI Core 内部有一块本地内存 Local Memory，因为靠近计算单元，所以它的带宽会非常高，相对的容量就会很小，比如一般是几百 K 到 1M。AI Core 内部的核心组件有三个计算单元，标量计算单元、向量计算单元，矩阵计算单元。另外还有一个 DMA 搬运单元，DMA 搬运单元负责在 Global Memory 和 Local Memory 之间搬运数据。
AI Core 内部的异步并行计算过程：Scalar 计算单元读取指令序列，并把向量计算、矩阵计算、数据搬运指令发射给对应单元的指令队列，向量计算单元、矩阵计算单元、数据搬运单元异步并行执行接收到的指令。该过程可以参考上图中蓝色箭头所示的指令流。不同的指令间有可能存在依赖关系，为了保证不同指令队列间的指令按照正确的逻辑关系执行，Scalar 计算单元也会给对应单元下发同步指令。各单元之间的同步过程可以参考上图中的橙色箭头所示的同步信号流。
AI Core 内部数据处理的基本过程：DMA 搬入单元把数据搬运到 Local Memory，Vector/Cube 计算单元完成数据，并把计算结果写回 Local Memory，DMA 搬出单元把处理好的数据搬运回 Global Memory。该过程可以参考上图中的红色箭头所示的数据流。
Ascend C 编程模型基础
Ascend C 编程范式
Ascend C 编程范式是一种流水线式的编程范式，把算子核内的处理程序，分成多个流水任务，通过队列（Queue）完成任务间通信和同步，并通过统一的内存管理模块（Pipe）管理任务间通信内存。流水编程范式应用了流水线并行计算方法。
若 n=3，即待处理的数据被切分成 3 片，则上图中的流水任务运行起来的示意图如下，从运行图中可以看出，对于同一片数据，Stage1、Stage2、Stage3 之间的处理具有依赖关系，需要串行处理；不同的数据切片，同一时间点，可以有多个任务在并行处理，由此达到任务并行、提升性能的目的。
Ascend C 分别针对 Vector、Cube 编程设计了不同的流水任务。开发者只需要完成基本任务的代码实现即可，底层的指令同步和并行调度由 Ascend C 框架实现，开发者无需关注。
矢量编程范式
矢量编程范式把算子的实现流程分为 3 个基本任务：CopyIn，Compute，CopyOut。CopyIn 负责搬入操作，Compute 负责矢量计算操作，CopyOut 负责搬出操作。
我们只需要根据编程范式完成基本任务的代码实现就可以了，底层的指令同步和并行调度由 Ascend C 框架来实现。
那 Ascend C 是怎么完成不同任务之间的数据通信和同步的呢？这里 Ascend C 提供了 Queue 队列管理的 API，主要就是两个队列操作 API EnQue、DeQue 以及内存的逻辑抽象。
矢量编程中使用到的逻辑位置（QuePosition）定义如下：
搬入数据的存放位置：VECIN；
计算中间变量的位置：VECCALC；
搬出数据的存放位置：VECOUT。
从前面可以看到，矢量编程主要分为 CopyIn、Compute、CopyOut 三个任务。CopyIn 任务中将输入数据从 Global 内存搬运至 Local 内存后，需要使用 EnQue 将 LocalTensor 放入 VECIN 的 Queue 中；Compute 任务等待 VECIN 的 Queue 中 LocalTensor 出队之后才可以完成矢量计算，计算完成后使用 EnQue 将计算结果 LocalTensor 放入到 VECOUT 的 Queue 中；CopyOut 任务等待 VECOUT 的 Queue 中 LocalTensor 出队，再将其拷贝到 Global 内存。这样 ，Queue 队列就完成了三个任务间的数据通信和同步。具体流程和流程图如下：
使用 DataCopy 接口将 GlobalTensor 数据拷贝到 LocalTensor。
使用 EnQue 接口将 LocalTensor 放入 VECIN 的 Queue 中。
使用 DeQue 接口从 VECIN 中取出 LocalTensor。
使用 Ascend C 接口完成矢量计算。
使用 EnQue 接口将计算结果 LocalTensor 放入到 VECOUT 的 Queue 中。
使用 DeQue 接口从 VECOUT 的 Queue 中去除 LocalTensor。
使用 DataCopy 接口将 LocalTensor 拷贝到 GlobalTensor 上。
这样我们的 kernel 实现代码就很清晰了。先初始化内存和队列，然后通过编程范式实现 CopyIn、Compute、CopyOut 三个 Stage 就可以了。
SPMD 并行编程-多核
最前面介绍昇腾 AI 处理器的时候，有介绍过 AI Core 是有多个的，那我们怎么把多个 AI Core 充分利用起来呢？常用的并行计算方法中，有一种 SPMD（Single-Program Multiple-Data）数据并行的方法，简单说就是将数据分片，每片数据经过完整的一个数据处理流程。这个就能和昇腾 AI 处理器的多核匹配上了，我们将数据分成多份，每份数据的处理运行在一个核上，这样每份数据并行处理完成，整个数据也就处理完了。Ascend C 是 SPMD（Single-Program Multiple-Data）编程，多个 AI Core 共享相同的指令代码，每个核上的运行实例唯一的区别是就是 block_idx（内置变量）不同，这样我们就可以通过 block_idx 来区分不同的核，只要对 Global Memory 上的数据地址进行切分偏移，就可以让每个核处理自己对应的那部分数据了。
算子被调用时，所有的计算核心都执行相同的实现代码，入口函数的入参也是相同的。每个核上处理的数据地址需要在起始地址上增加 block_idx*BLOCK_LENGTH（每个 block 处理的数据长度）的偏移来获取。这样也就实现了多核并行计算的数据切分。
class KernelAdd {
public:
__aicore__ inline KernelAdd() {}
__aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z)
{
// get start index for current core, core parallel
GM_ADDR xGmOffset = x + BLOCK_LENGTH * GetBlockIdx();
GM_ADDR yGmOffset = y + BLOCK_LENGTH * GetBlockIdx();
GM_ADDR zGmOffset = z + BLOCK_LENGTH * GetBlockIdx();
xGm.SetGlobalBuffer((__gm__ half*)xGmOffset, BLOCK_LENGTH);
yGm.SetGlobalBuffer((__gm__ half*)yGmOffset, BLOCK_LENGTH);
zGm.SetGlobalBuffer((__gm__ half*)zGmOffset, BLOCK_LENGTH);
……
}
……
}
复制代码
Ascend C API 介绍
在整个 kernel 实现中，最最核心的代码就是 Add(zLocal, xLocal, yLocal, TILE_LENGTH);通过一个 Ascend C 提供的 API 接口完成了所有数据的加法计算，对，没看错，就是这个接口完成了计算。
接下来就介绍下 Ascend C 提供的 API。Ascend C 算子采用标准 C++语法和一组类库 API 进行编程，类库 API 主要包含以下几种，大家可以在核函数的实现中根据自己的需求选择合适的 API：
计算类 API，包括标量计算 API、向量计算 API、矩阵计算 API，分别实现调用 Scalar 计算单元、Vector 计算单元、Cube 计算单元执行计算的功能。
数据搬运 API，上述计算 API 基于 Local Memory 数据进行计算，所以数据需要先从 Global Memory 搬运至 Local Memory，再使用计算接口完成计算，最后从 Local Memory 搬出至 Global Memory。执行搬运过程的接口称之为数据搬移接口，比如 DataCopy 接口。
内存管理 API，用于分配管理内存，比如 AllocTensor、FreeTensor 接口。
任务同步 API，完成任务间的通信和同步，比如 EnQue、DeQue 接口。
Ascend C API 的计算操作数都是 Tensor 类型：GlobalTensor 和 LocalTensor。
介绍完 Ascend C API 种类后，下面来解释下为什么一个 Add 接口就可以计算所有的数。原来 Ascend C 编程模型是基于 SIMD（单指令多数据）架构的，单条指令可以完成多个数据操作，同时在 API 内部封装了一些指令的高级功能。
算子执行基本流程
前面有提到，在异构计算架构中，NPU 与 CPU 是协同工作的，在 Ascend C 编程模型中，我们需要实现 NPU 侧的代码和 CPU 侧的代码。在 NPU 侧的代码我们通常叫做 Kernel 实现代码，CPU 侧的代码我们一般叫做 Host 实现代码，一份完整的 Ascend C 代码，通常包括 Host 侧实现代码和 Kernel 侧实现代码。Ascend C 算子执行的基本流程如下：
初始化 Device 设备；
创建 Context 绑定设备；
分配 Host 内存，并进行数据初始化；
分配 Device 内存，并将数据从 Host 上拷贝到 Device 上；
用内核调用符<<<>>>调用核函数完成指定的运算；
将 Device 上的运算结果拷贝回 Host；
释放申请的资源。
核函数介绍
上面的流程中，最重要的一步就是调用核函数来进行并行计算任务。核函数（Kernel Function）是 Ascend C 算子 Device 侧实现的入口。在核函数中，需要为在 AI 核上执行的代码规定要进行的数据访问和计算操作。
extern "C" __global__ __aicore__ void add_custom(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z);
复制代码
上面这个是一个核函数声明的示例，extern "C"表示核函数按照类 C 的编译和连接规约来编译和连接，__global__函数类型限定符表示它是一个核函数， __aicore__函数类型限定符表示该核函数在 device 侧的 AI Core 上执行。参数列表中的变量类型限定符__gm__，表明该指针变量指向 Global Memory 上某处内存地址，注意这里的入参只能支持指针或 C/C++内置数据类型，样例里指针使用的类型为 uint8_t，在后续的使用中需要将其转化为实际的指针类型。
Ascend C 编程模型中的核函数采用内核调用符<<<...>>>来调用，样例如下：
kernel_name<<<blockDim, l2ctrl, stream>>>(argument list);
复制代码
kernel_name 即为上面讲的核函数名称，argument list 是核函数的函数入参，在<<<>>>中间，有 3 个参数：
blockDim，规定了核函数将会在几个核上执行，我们可以先设置为 1；
l2ctrl，保留参数，暂时设置为固定值 nullptr，我们不用关注；
stream，使用 aclrtCreateStream 创建，用于多线程调度。
样例开发讲解
样例代码结构
|-- CMakeLists.txt //编译工程文件
|-- cmake //编译工程文件
|-- data_utils.h //数据读入写出函数
|-- input //存放脚本生成的输入数据目录
|-- leakyrelu_custom.cpp //算子kernel实现
|-- leakyrelu_custom.py //输入数据和真值数据生成脚本文件
|-- leakyrelu_custom_tiling.h //host侧tiling函数
|-- main.cpp //主函数，host侧调用代码，含cpu域及npu域调用
|-- output //存放算子运行输出数据和标杆数据的目录
|-- readme.md //执行命令说明
|-- run.sh //运行脚本
复制代码
主要文件
输入数据和真值数据生成脚本文件：KERNEL_NAME.py。
根据算子的输入输出编写生成输入数据和真值数据的脚本。
本例子生成 8 * 200 * 1024 大小的 fp16 数据：
……
def gen_golden_data_simple():
total_length_imm = 8 * 200 * 1024
tile_num_imm = 8
//生成tilling的bin文件
total_length = np.array(total_length_imm, dtype=np.uint32)
tile_num = np.array(tile_num_imm, dtype=np.uint32)
scalar = np.array(0.1, dtype=np.float32)
tiling = (total_length, tile_num, scalar)
tiling_data = b''.join(x.tobytes() for x in tiling)
with os.fdopen(os.open('./input/tiling.bin', WRITE_FILE_FLAGS, PEN_FILE_MODES_640), 'wb') as f:
f.write(tiling_data)
//生成输入数据
input_x = np.random.uniform(-100, 100, [8, 200, 1024]).astype(np.float16)
//生成golden数据，功能和LeakyRelu相同
golden = np.where(input_x > 0, input_x, input_x * scalar).astype(np.float16)
input_x.tofile("./input/input_x.bin")
golden.tofile("./output/golden.bin")
复制代码
编译工程文件：CMakeLists.txt
用于编译 cpu 侧或 npu 侧运行的 Ascend C 算子。主要关注 CMakeLists.txt 中源文件是否全部列全。
调用算子的应用程序：main.cpp
主要是内存申请，数据拷贝和文件读写等操作，并最终调用算子，相关 API 的介绍如下：
1.AscendCL 初始化接口 aclInit，用于运行时接口 AscendCL 的初始化，是程序最先调用的接口；aclrtCreateContext 和 aclrtCreateStream 用于创建 Context 和 Stream，主要用于线程相关的资源管理。
2.aclrtMallocHost 接口，用于在 Host 上申请内存：
aclError aclrtMallocHost(void **hostPtr, size_t size)
这个函数和 C 语言中的 malloc 类似，用于在 Host 上申请一定字节大小的内存，其中 hostPtr 是指向所分配内存的指针，size 是申请的内存大小，如果需要释放这块内存的话，使用 aclrtFreeHost 接口释放，这和 C 语言中的 free 函数对应。
3.aclrtMalloc 接口，用于在 Device 上申请内存：
aclError aclrtMalloc(void **devPtr, size_t size, aclrtMemMallocPolicy policy)
和 Host 上的内存申请接口相比，多了一个 policy 参数，用于设置内存分配规则，一般设置成 ACL_MEM_MALLOC_HUGE_FIRST 就可以了。使用完毕后可以用对应的aclrtFree接口释放内存。
4.aclrtMemcpy 接口，用于 Host 和 Device 之间数据拷贝：
前面申请的内存区分了 Host 内存和 Device 内存，那就会涉及到数据同步的问题，aclrtMemcpy 就是用于 Host 和 Device 之间数据通信的接口：
aclError aclrtMemcpy(void *dst, size_t destMax, const void *src, size_t count, aclrtMemcpyKind kind)
其中 src 指向数据源，而 dst 是目标内存地址，destMax 是目的内存地址的最大内存长度，count 是拷贝的字节数，其中 aclrtMemcpyKind 控制复制的方向：ACL_MEMCPY_HOST_TO_HOST、ACL_MEMCPY_HOST_TO_DEVICE、ACL_MEMCPY_DEVICE_TO_HOST 和 ACL_MEMCPY_DEVICE_TO_DEVICE，像 ACL_MEMCPY_HOST_TO_DEVICE 就是将 Host 上数据拷贝到 Device 上。
5.核心函数为 CPU 侧的调用 kernel 函数
ICPU_RUN_KF(leakyrelu_custom, blockDim, x, y, usrWorkSpace, tiling);
复制代码
和 NPU 侧调用的
leakyrelu_custom_do(blockDim, nullptr, stream, xDevice, yDevice, workspaceDevice, tilingDevice);
复制代码
完整代码如下：
//This file constains code of cpu debug and npu code.We read data from bin file and write result to file.
#include "data_utils.h"
#include "leakyrelu_custom_tiling.h"
#ifndef __CCE_KT_TEST__
#include "acl/acl.h"
extern void leakyrelu_custom_do(uint32_t coreDim, void* l2ctrl, void* stream, uint8_t* x, uint8_t* y,
uint8_t* workspace, uint8_t* tiling);
#else
#include "tikicpulib.h"
extern "C" __global__ __aicore__ void leakyrelu_custom(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling);
#endif
int32_t main(int32_t argc, char* argv[])
{
size_t tilingSize = sizeof(LeakyReluCustomTilingData);
size_t usrWorkspaceSize = 4096;
size_t sysWorkspaceSize = 16 * 1024 * 1024;
uint32_t blockDim = 8;
#ifdef __CCE_KT_TEST__ //CPU侧调用
//申请内存用于存放workspace和tilling数据
uint8_t* usrWorkSpace = (uint8_t*)AscendC::GmAlloc(usrWorkspaceSize);
uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(tilingSize);
ReadFile("./input/tiling.bin", tilingSize, tiling, tilingSize);
size_t inputByteSize = blockDim * 200 * 1024 * sizeof(uint16_t); // uint16_t represent half
size_t outputByteSize = blockDim * 200 * 1024 * sizeof(uint16_t); // uint16_t represent half
//申请内存用于存放输入和输出数据
uint8_t* x = (uint8_t*)AscendC::GmAlloc(inputByteSize);
uint8_t* y = (uint8_t*)AscendC::GmAlloc(inputByteSize);
//获取输入数据
ReadFile("./input/input_x.bin", inputByteSize, x, inputByteSize);
// PrintData(x, 16, printDataType::HALF);
//在AIV上执行
AscendC::SetKernelMode(KernelMode::AIV_MODE);
//调用kernel函数
ICPU_RUN_KF(leakyrelu_custom, blockDim, x, y, usrWorkSpace, tiling); // use this macro for cpu debug
// PrintData(y, 16, printDataType::HALF);
WriteFile("./output/output_y.bin", y, outputByteSize);
AscendC::GmFree((void *)x);
AscendC::GmFree((void *)y);
AscendC::GmFree((void *)usrWorkSpace);
AscendC::GmFree((void *)tiling);
#else //NPU侧调用
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));
uint8_t *xHost, *yHost, *tilingHost, *workspaceHost;
uint8_t *xDevice, *yDevice, *tilingDevice, *workspaceDevice;
//申请host上tilling内存并读入tilling数据
CHECK_ACL(aclrtMallocHost((void**)(&tilingHost), tilingSize));
ReadFile("./input/tiling.bin", tilingSize, tilingHost, tilingSize);
//申请host上workspace内存
CHECK_ACL(aclrtMallocHost((void**)(&workspaceHost), tilingSize));
size_t inputByteSize = blockDim * 200 * 1024 * sizeof(uint16_t); // uint16_t represent half
size_t outputByteSize = blockDim * 200 * 1024 * sizeof(uint16_t); // uint16_t represent half
size_t workspaceByteSize = sysWorkspaceSize + usrWorkspaceSize;
//申请host和device上的输入输出内存和device上的workspace和tilling内存
CHECK_ACL(aclrtMallocHost((void**)(&xHost), inputByteSize));
CHECK_ACL(aclrtMallocHost((void**)(&yHost), inputByteSize));
CHECK_ACL(aclrtMallocHost((void**)(&workspaceHost), workspaceByteSize));
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**)&tilingDevice, tilingSize, ACL_MEM_MALLOC_HUGE_FIRST));
CHECK_ACL(aclrtMalloc((void**)&workspaceDevice, workspaceByteSize, ACL_MEM_MALLOC_HUGE_FIRST));
ReadFile("./input/input_x.bin", inputByteSize, xHost, inputByteSize);
// PrintData(xHost, 16, printDataType::HALF);
//从host上拷贝输入数据和tilling数据到device
CHECK_ACL(aclrtMemcpy(xDevice, inputByteSize, xHost, inputByteSize, ACL_MEMCPY_HOST_TO_DEVICE));
CHECK_ACL(aclrtMemcpy(tilingDevice, tilingSize, tilingHost, tilingSize, ACL_MEMCPY_HOST_TO_DEVICE));
//调用核函数
leakyrelu_custom_do(blockDim, nullptr, stream, xDevice, yDevice, workspaceDevice, tilingDevice);
//等待核函数运行完成
CHECK_ACL(aclrtSynchronizeStream(stream));
//拷回运行结果到host
CHECK_ACL(aclrtMemcpy(yHost, outputByteSize, yDevice, outputByteSize, ACL_MEMCPY_DEVICE_TO_HOST));
// PrintData(yHost, 16, printDataType::HALF);
WriteFile("./output/output_y.bin", yHost, outputByteSize);
//释放资源
CHECK_ACL(aclrtFree(xDevice));
CHECK_ACL(aclrtFree(yDevice));
CHECK_ACL(aclrtFree(workspaceDevice));
CHECK_ACL(aclrtFree(tilingDevice));
CHECK_ACL(aclrtFreeHost(xHost));
CHECK_ACL(aclrtFreeHost(yHost));
CHECK_ACL(aclrtFreeHost(workspaceHost));
CHECK_ACL(aclrtFreeHost(tilingHost));
CHECK_ACL(aclrtDestroyStream(stream));
CHECK_ACL(aclrtDestroyContext(context));
CHECK_ACL(aclrtResetDevice(deviceId));
CHECK_ACL(aclFinalize());
#endif
return 0;
}
复制代码
一键式编译运行脚本 run.sh
编译和运行应用程序。
cpu 侧运行命令：
bash run.sh leakyrelu_custom ascend910B1 VectorCore cpu
复制代码
npu 侧运行命令：
bash run.sh leakyrelu_custom ascend910B1 VectorCore npu
复制代码
参数含义如下：
bash run.sh <kernel_name> <soc_version> <core_type> <run_mode>
复制代码
<kernel_name>表示需要运行的算子。
<soc_version>表示算子运行的 AI 处理器型号。
<core_type>表示在 AI Core 上或者 Vector Core 上运行，参数取值为 AiCore/VectorCore。
<run_mode>表示算子以 cpu 模式或 npu 模式运行，参数取值为 cpu/npu。
kernel 实现
函数原型定义
本样例中，函数名为 leakyrelu_custom，根据对算子输入输出的分析，确定有 2 个参数 x，y，其中 x 为输入内存，y 为输出内存。核函数原型定义如下所示：
extern "C" __global__ __aicore__ void leakyrelu_custom(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling){ }
复制代码
使用__global__函数类型限定符来标识它是一个核函数，可以被<<<...>>>调用；使用__aicore__函数类型限定符来标识该核函数在设备端 AI Core 上执行；为方便起见，统一使用 GM_ADDR 宏修饰入参，GM_ADDR 宏定义：
#define GM_ADDR __gm__ uint8_t* __restrict__
复制代码
获取 tilling 数据，并调用算子类的 Init 和 Process 函数。
算子类的 Init 函数，完成内存初始化相关工作，Process 函数完成算子实现的核心逻辑。
extern "C" __global__ __aicore__ void leakyrelu_custom(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
{
GET_TILING_DATA(tilingData, tiling);
KernelLeakyRelu op;
op.Init(x, y, tilingData.totalLength, tilingData.tileNum, tilingData.scalar);
op.Process();
}
复制代码
对核函数的调用进行封装
封装后得到 leakyrelu_custom_do 函数，便于主程序调用。#ifndef __CCE_KT_TEST__表示该封装函数仅在编译运行 NPU 侧的算子时会用到，编译运行 CPU 侧的算子时，可以直接调用 add_custom 函数。调用核函数时，除了需要传入输入输出参数 x，y，切分相关参数 tiling，还需要传入 blockDim（核函数执行的核数）, l2ctrl（保留参数，设置为 nullptr）, stream（应用程序中维护异步操作执行顺序的 stream）来规定核函数的执行配置。
#ifndef __CCE_KT_TEST__
// call of kernel function
void leakyrelu_custom_do(uint32_t blockDim, void* l2ctrl, void* stream, uint8_t* x, uint8_t* y,
uint8_t* workspace, uint8_t* tiling)
{
leakyrelu_custom<<<blockDim, l2ctrl, stream>>>(x, y, workspace, tiling);
}
#endif
复制代码
获取 tiling 参数
主要从 tilingPointer 中获取 tiling 的参数 totalLength(总长度)、tileNum（切分个数，单核循环处理数据次数）和 scalar（LeakyRelu 计算标量）。
#define GET_TILING_DATA(tilingData, tilingPointer) \
LeakyReluCustomTilingData tilingData; \
INIT_TILING_DATA(LeakyReluCustomTilingData, tilingDataPointer, tilingPointer); \
(tilingData).totalLength = tilingDataPointer->totalLength; \
(tilingData).tileNum = tilingDataPointer->tileNum; \
(tilingData).scalar = tilingDataPointer->scalar;
#endif // LEAKYRELU_CUSTOM_TILING_H
复制代码
Init 函数
主要获取 tiling 数据后，设置单核上 gm 的地址和 Buffer 的初始化。
__aicore__ inline void Init(GM_ADDR x, GM_ADDR y, uint32_t totalLength, uint32_t tileNum, float scalar)
{
ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");
this->blockLength = totalLength / GetBlockNum();
this->tileNum = tileNum;
this->scalar = static_cast<half>(scalar);
ASSERT(tileNum != 0 && "tile num can not be zero!");
this->tileLength = this->blockLength / tileNum / BUFFER_NUM;
// get start index for current core, core parallel
xGm.SetGlobalBuffer((__gm__ half*)x + this->blockLength * get_block_idx(), this->blockLength);
yGm.SetGlobalBuffer((__gm__ half*)y + this->blockLength * get_block_idx(), this->blockLength);
// pipe alloc memory to queue, the unit is Bytes
pipe.InitBuffer(inQueueX, BUFFER_NUM, this->tileLength * sizeof(half));
pipe.InitBuffer(outQueueY, BUFFER_NUM, this->tileLength * sizeof(half));
}
复制代码
Process 函数
主要实现三个 CopyIn、Compute、CopyOut 这三 stage。
__aicore__ inline void Process()
{
// loop count need to be doubled, due to double buffer
int32_t loopCount = this->tileNum * BUFFER_NUM;
// tiling strategy, pipeline parallel
for (int32_t i = 0; i < loopCount; i++) {
CopyIn(i);
Compute(i);
CopyOut(i);
}
}
复制代码
CopyIn 函数
负责从 Global Memory 拷贝数据到 Local Memory，并将数据加入 Queue
__aicore__ inline void CopyIn(int32_t progress)
{
// alloc tensor from queue memory
LocalTensor<half> xLocal = inQueueX.AllocTensor<half>();
// copy progress_th tile from global tensor to local tensor
DataCopy(xLocal, xGm[progress * tileLength], tileLength);
// enque input tensors to VECIN queue
inQueueX.EnQue(xLocal);
}
复制代码
Compute 函数
负责从 Queue 中取出数据，进行计算，并将结果放入 Queue
__aicore__ inline void Compute(int32_t progress)
{
// deque input tensors from VECIN queue
LocalTensor<half> xLocal = inQueueX.DeQue<half>();
LocalTensor<half> yLocal = outQueueY.AllocTensor<half>();
// call LeakyRelu instr for computation
LeakyRelu(yLocal, xLocal, scalar, tileLength);
// enque the output tensor to VECOUT queue
outQueueY.EnQue<half>(yLocal);
// free input tensors for reuse
inQueueX.FreeTensor(xLocal);
}
复制代码
CopyOut 函数
负责从 Queue 中将数据取出，并将数据从 Local Memory 拷贝到 Global Memory。
__aicore__ inline void CopyOut(int32_t progress)
{
// deque output tensor from VECOUT queue
LocalTensor<half> yLocal = outQueueY.DeQue<half>();
// copy progress_th tile from local tensor to global tensor
DataCopy(yGm[progress * tileLength], yLocal, tileLength);
// free output tensor for reuse
outQueueY.FreeTensor(yLocal);
}
复制代码
编译和执行
在 CPU 侧执行
执行结果如下：
可以看到最后的输出结果 output_y.bin 和标杆数据 golden.bin 的 MD5 值相同，说明计算结果相同。
执行完成后，在 input 下存放输入数据和 tiling 数据，在 output 下面存放了输出数据和标杆数据，npuchk 目录下是每个核的 npu_check 执行结果。
在当前目录还有一个可执行二进制文件 leakyrelu_custom_cpu，如果执行报错，可以通过 gdb 调试这个可执行文件，具体调试可参考文末官方教程。
在 NPU 侧执行
在 NPU 侧执行有两种方式：仿真执行和上板运行，命令都相同，只是编译选项不同，我们可以通过修改编译选项-DASCEND_RUN_MODE 为 SIMULATOR 运行 CAModel 仿真，设置为 ONBOARD 是上板运行。
function compile_and_execute() {
# 使用cmake编译cpu侧或者npu侧算子, SIMULATOR or ONBOARD
mkdir -p build; cd build; \
cmake .. \
-Dsmoke_testcase=$1 \
-DASCEND_PRODUCT_TYPE=$2 \
-DASCEND_CORE_TYPE=$3 \
-DASCEND_RUN_MODE="SIMULATOR" \
-DASCEND_INSTALL_PATH=$ASCEND_HOME_DIR
VERBOSE=1 cmake --build . --target ${1}_${4}
……
}
复制代码
参考资料
总之，学习 Ascend C，仅需了解 C++编程、理解对列通信与内存申请释放机制、通过调用相应的计算接口与搬运接口，就可以写出运行在昇腾 AI 处理器上的高性能算子。
了解更多 Ascend C 学习资源，请访问官方教程：Ascend C编程指南（官方教程）
号外！
华为将于 2023 年 9 月 20-22 日，在上海世博展览馆和上海世博中心举办第八届华为全联接大会（HUAWEICONNECT 2023）。本次大会以“加速行业智能化”为主题，邀请思想领袖、商业精英、技术专家、合作伙伴、开发者等业界同仁，从商业、产业、生态等方面探讨如何加速行业智能化。
我们诚邀您莅临现场，分享智能化的机遇和挑战，共商智能化的关键举措，体验智能化技术的创新和应用。您可以：
在 100+场主题演讲、峰会、论坛中，碰撞加速行业智能化的观点
参观 17000 平米展区，近距离感受智能化技术在行业中的创新和应用
与技术专家面对面交流，了解最新的解决方案、开发工具并动手实践
与客户和伙伴共寻商机
感谢您一如既往的支持和信赖，我们热忱期待与您在上海见面。
大会官网：https://www.huawei.com/cn/events/huaweiconnect
欢迎关注“华为云开发者联盟”公众号，获取大会议程、精彩活动和前沿干货。
点击关注，第一时间了解华为云新鲜技术~
评论