1 简介
相信AI开发者对CUDA编程平台再也熟悉不过了。几乎每一个AI初学者第一个接触的硬件加速都是NVIDIA CUDA。NVIDIA有廉价的游戏卡提供高性能高性价比的AI算力，所以是学生时期最容易入手的AI加速廉价硬件。
然而，华为的昇腾平台主要面向政府和大型企业，对于学生或小型企业来说价格高昂，很少在零售平台出现，所以接触较少。
因此，本文详细介绍下昇腾Ascend C编程与NVIDIA CUDA编程的差异。同时也为两者之间的程序迁移提供初步思路。同时，让读者更熟悉AscendC的开发环境op.Process();。
2 主机代码
2.1上下文和流
Ascend C和CUDA一样同样也需要主机C代码对上下文和设置内存进行操控。不过，Ascend C的上下文是需要显式操作的，并不像CUDA是隐式自动的。
例如：
AscendC的上下文初始化
#pragma once
#include <acl/acl.h>
#define CHECK_ACL(x) \
do { \
aclError __ret = x; \
if (__ret != ACL_ERROR_NONE) { \
std::cerr << __FILE__ << ":" << __LINE__ << " aclError:" << __ret << std::endl; \
} \
} while (0);
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));
CHECK_ACL是检验aclrt接口返回的错误值，这一点和CUDA一样，返回值为0，则是接口调用成功。
CUDA对于此操作是可选的，可以不写，自动创建。 如果需要创建的额外的显式流，调用cudaCreateStream即可。
2.2 设备内存
设备内存的接口函数，两者都是对称的。
| 作用 | AscendC | CUDA |
|---|---|---|
| 分配 | aclrtMalloc | cudaMalloc |
| 拷贝 | aclrtMemcpy | cudaMemcpy |
| 释放 | aclrtFree | cudaFree |
AscendC示范：
CHECK_ACL(aclrtMalloc((void **)&d_a, m * n * sizeof(float),
ACL_MEM_MALLOC_HUGE_FIRST));
CHECK_ACL(aclrtMemcpy(d_a, cpu_a.size() * sizeof(float), cpu_a.data(),
cpu_a.size() * sizeof(float),
ACL_MEMCPY_HOST_TO_DEVICE));
CUDA示范：
CHECK_CUDA(cudaMalloc((void **)&d_a, m * n * sizeof(float)));
CHECK_CUDA(aclrtMemcpy(d_a, cpu_a.data(),
cpu_a.size() * sizeof(float),
cudaMemcpyHostToDevice));
aclrtMalloc在分配上了多了一个参数，内存分配策略。默认使用ACL_MEM_MALLOC_HUGE_FIRST，这是分配页内存的方式。优先试用普通页，超过1MB时，分配大页内存。
aclrtMemcpy比cudaMemcpy少写一个拷贝大小的参数。提供了一个目标大小限制的安全范围。
2.3 启动Kernel
昇腾启动只需要填写一个blockDim和函数对应的参数即可。
ACLRT_LAUNCH_KERNEL(polar_kernel)
(BLOCK_DIM, stream, d_a, d_b, d_c, workspaceDevice, tilingDevice);
NVIDIA需要两个参数GRID_DIM和BLOCK_DIM。
polar_kernel <<< GRID_DIM， BLOCK_DIM >>>(d_a, d_b, d_c);
昇腾硬件是SPMD模型。BLOCK_DIM是指物理核启动的数量，如910b一共只有40个vector core，那么实际上最多启动40个。
而CUDA是SIMT模型，包含大量了线程束。一个GRID由大量的BLOCK组成xyz阵列，就是GRID_DIM。一个BLOCK由大量的xyz阵列的线程Thread组成，就是BLOCK_DIM。
昇腾的线程束很少，一般是20-40个。而NVIDIA的线程数可能由几千或几万个。因此，编写kernel时，不能采用相同的并行算法。
3 设备代码
3.1 昇腾的kernel加法kernel最简单实现示范
template <typename DTYPE_X>
class KernelAdd
{
public:
__aicore__ inline KernelPolar() {}
__aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, int64_t length)
{
this->loopCount = 16;
this->tileLength = length / GetBlockNum() / loopCount;
pipe.InitBuffer(inQueueX, BUFFER_NUM, this->tileLength * sizeof(DTYPE_X));
pipe.InitBuffer(inQueueY, BUFFER_NUM, this->tileLength * sizeof(DTYPE_X));
pipe.InitBuffer(outQueueZ, BUFFER_NUM, this->tileLength * sizeof(DTYPE_X));
}
__aicore__ inline void Process()
{
for (int i = 0; i < this->loopCount; i++)
{
CopyIn(i);
Compute(i);
CopyOut(i);
}
}
private:
__aicore__ inline void CopyIn(int index)
{
LocalTensor<DTYPE_X> xLocal = inQueueX.AllocTensor<DTYPE_X>();
LocalTensor<DTYPE_X> yLocal = inQueueY.AllocTensor<DTYPE_X>();
DataCopy(xLocal, xGm[index * this->tileLength], this->tileLength);
DataCopy(yLocal, yGm[index * this->tileLength], this->tileLength);
inQueueX.EnQue(xLocal);
inQueueY.EnQue(yLocal);
}
__aicore__ inline void Compute(int index)
{
LocalTensor<DTYPE_X> xLocal = inQueueX.DeQue<DTYPE_X>();
LocalTensor<DTYPE_X> yLocal = inQueueY.DeQue<DTYPE_X>();
LocalTensor<DTYPE_X> zLocal = outQueueZ.AllocTensor<DTYPE_X>();
Add(zLocal, xLocal, yLocal, this->tileLength);
outQueueZ.EnQue<DTYPE_X>(zLocal);
inQueueX.FreeTensor(xLocal);
inQueueY.FreeTensor(yLocal);
}
__aicore__ inline void CopyOut(int index)
{
LocalTensor<DTYPE_X> zLocal = outQueueZ.DeQue<DTYPE_X>();
DataCopy(zGm[index * this->tileLength], zLocal, this->tileLength);
outQueueZ.FreeTensor(zLocal);
}
private:
TPipe pipe;
TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX, inQueueY;
TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueZ;
GlobalTensor<DTYPE_X> xGm;
GlobalTensor<DTYPE_X> yGm;
GlobalTensor<DTYPE_X> zGm;
uint32_t tileLength{};
uint32_t loopCount = 0;
};
昇腾的kernel入口函数必须声明extern "C" __global__ __aicore__ 。必须是C语言的函数符号，并且声明__aicore__表示在aicore上运算。
参数习惯放在tiling结构体中。通过GET_TILING_DATA(tiling_data, tiling);将数据取出到寄存器上。模板类型使用编译器内置的DTYPE_X，这是根据算子定义生成的，“DTYPE_输入参数大写”。
昇腾的kernel编写更倾向于将任务平均分配到每个block上。然后每个block拆分成若干个串行的loop，执行双缓冲流水。
op.Init(x, y, z, tiling_data.alpha，tiling_data.length); 实现对参数的传参。以及pipe.InitBuffer对全局内存向GlobalTensor绑定，设置每个Block的可访问的内存区间。
op.Process(); 是实现计算双缓冲流水的三部曲封装：CopyIn拷入、Compute运算、CopyOut拷出。
昇腾的计算必须要在LocalTensor上进行，CopyIn通过DataCopy将GlobalTensor复制到LocalTensor。Process执行真正的运算Add(zLocal, xLocal, yLocal, this->tileLength); 一次计算tileLength个矢量加法。CopyOut通过DataCopy将LocalTensor复制到GLobalTensor上，完成一次Loop的计算。待所有Loop完成后，当前block的任务结束。
3.2 CUDA加法最简单Kernel实现示范
__global__ void custom_add(float* x, float* y, float* z, float alpha, size_t N) {
int index = threadIdx.x + blockDim.x * blockIdx.x;
if (index < N)
{
z[index] = x[index] + y[index];
}
}
CUDA的核函数比较简单。SIMT的编程模型，只需要将加法分配到每个线程上即可，通过int index = threadIdx.x + blockDim.x * blockIdx.x; 计算出当前线程ID。
不需要手动管理全局内存和局部内存之间的转换，或者说全局内存的数组可以直接使用运算符进行计算，如**z[index] = x[index] + y[index]**。
4 总结
昇腾AscendC与CUDA之间的代码编写差异是非常大的。昇腾对核函数需要的细节更多，需要手动管理全局内存(GlobalTensor)的可见区间，和局部内存(LocalTensor)在片上缓存的手动控制。不像CUDA那样，可以直接调用全局内存运算，忽略全局到局部的细节处理。同时，昇腾属于SPMD的NPU结构，需要算子开发者对昇腾的硬件架构有深刻的理解，和算子算法的特有的优化方法。
