Ascend C 学习笔记

AI 处理器#

组成#

HBM（存储）
DVPP（硬件编解码）
CPU（4 核，ARM）：指令分发等功能
AI Core NPU（AI 计算）：执行 AI 计算，矩阵计算（16x16x16），标量计算，向量计算（128）等。

通用计算与 AI 计算的区别#

通用计算：适用于大多数的计算任务，比如加减乘除等。
比如计算矩阵乘法，只能串行计算，速度慢。

NPU 计算，可以使用向量计算单元，同时计算多个数据，速度更快，或者直接使用矩阵计算单元，直接计算矩阵乘法。

异构计算：CPU + NPU，CPU 负责控制，NPU 负责计算。

Ascend C#

原生算子开发，自动调度。

图优化，自动流水。
算子深度融合，多个不同的算子，自动融合，减少数据内存拷贝

计算图执行下沉，将所有算子都下沉到 NPU 上执行。

自动调优 AOE

Ascend C#

AI Core#

计算单元，存储单元，控制单元。

20240711102909

内部结构：
标量计算单元（CPU）下发指令
20240711103146

Tensor 概念
name
shape
dtype
format

format#

N： batch size，图像数量
H： Height，高度
W： Width，宽度
C： Channel，通道数，RGB 为 3

Ascend C 优势#

20240711105258

Device 模块#

Host 与 Device

Host 指服务器
Device 指 NPU

Ascend C 算子开发#

Add 算子开发#

算子注册#

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();
}

在这里定义了我们这个 kernel 算子函数 add_custom，这个函数接受三个参数，分别是 x，y，z，这三个参数分别是输入的两个张量和输出的张量。

GM_ADDR 是 __gm__ uint8_t* 的别名。表示 global memory 上的地址。

算子实现#

在这里使用了一个类来实现算子。

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:
    ...
};

类中 public 了两个主要的函数，Init 和 Process。

__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));
}

Init 主要是对内存地址进行分配以及初始化。

__aicore__ inline void Process()
{
    int32_t loopCount = TILE_NUM * BUFFER_NUM;
    for (int32_t i = 0; i < loopCount; i++) {
        CopyIn(i);
        Compute(i);
        CopyOut(i);
    }
}

Process 函数循环调用了 CopyIn、Compute、CopyOut 这三个组成流水的函数。

其中 CopyIn 函数是将数据从全局内存拷贝到局部内存，Compute 函数是计算，CopyOut 函数是将计算结果从局部内存拷贝到全局内存。

__aicore__ inline void CopyIn(int32_t progress)
{
    LocalTensor<half> xLocal = inQueueX.AllocTensor<half>();
    LocalTensor<half> yLocal = inQueueY.AllocTensor<half>();
    DataCopy(xLocal, xGm[progress * TILE_LENGTH], TILE_LENGTH);
    DataCopy(yLocal, yGm[progress * TILE_LENGTH], TILE_LENGTH);
    inQueueX.EnQue(xLocal);
    inQueueY.EnQue(yLocal);
}
__aicore__ inline void Compute(int32_t progress)
{
    LocalTensor<half> xLocal = inQueueX.DeQue<half>();
    LocalTensor<half> yLocal = inQueueY.DeQue<half>();
    LocalTensor<half> zLocal = outQueueZ.AllocTensor<half>();
    Add(zLocal, xLocal, yLocal, TILE_LENGTH);
    outQueueZ.EnQue<half>(zLocal);
    inQueueX.FreeTensor(xLocal);
    inQueueY.FreeTensor(yLocal);
}
__aicore__ inline void CopyOut(int32_t progress)
{
    LocalTensor<half> zLocal = outQueueZ.DeQue<half>();
    DataCopy(zGm[progress * TILE_LENGTH], zLocal, TILE_LENGTH);
    outQueueZ.FreeTensor(zLocal);
}

在计算中，使用 Add 向量加法来实现这个算子。

由于所有的指令都会并发执行，所以这个算子中使用了 queue 来保证数据依赖。

DeQue 会阻塞，直到队列中有需要的数据。

算子调用#

在调用内核算子时

需要使用 <<<...>>> 来对函数进行调用。

blockDim：根据不同的设备，设置不同的 block 大小。
l2ctrl：目前为 nullptr
stream：表示整体的流

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);
}

Sinh 算子开发#

Sinh 的其他的结构与 Add 类似，区别主要是参数数量的差异。

其核心的区别点在于 Compute 函数的实现。

这里我使用了 Muls Exp，先求出了 $e^{-x}$
然后求出了 $e^x$
最后相减，再除以 2，就得到了 sinh 的结果。

__aicore__ inline void Compute(int32_t progress)
{
    LocalTensor<DTYPE_X> xLocal = inQueueX.DeQue<DTYPE_X>();
    LocalTensor<DTYPE_Y> yLocal = outQueueY.AllocTensor<DTYPE_Y>();
    LocalTensor<DTYPE_Y> tmpLocal = tmpBuffer.AllocTensor<DTYPE_Y>();
    // e^-x
    Muls(tmpLocal, xLocal, static_cast<DTYPE_X>(-1), this->tileLength);
    Exp(tmpLocal, tmpLocal, this->tileLength);
    // e^x 
    Exp(yLocal, xLocal, this->tileLength);
    // (e^x - e^-x)
    Sub(yLocal, yLocal, tmpLocal, this->tileLength);
    // (e^x - e^-x) / 2
    Muls(yLocal, yLocal, static_cast<DTYPE_X>(0.5), this->tileLength);
    outQueueY.EnQue<DTYPE_Y>(yLocal);
    inQueueX.FreeTensor(xLocal);
    tmpBuffer.FreeTensor(tmpLocal);
}