昇腾社区首页
支持
消息
我的论坛主页
论坛首页
/
【CANN训练营快速上手实践操作记录】Ascend C 算子编程范式及算子实现
收藏回复举报
【CANN训练营快速上手实践操作记录】Ascend C 算子编程范式及算子实现
发表于2023-05-31 21:34:41
0 查看

编程范式

编程范式描述了算子实现的固定流程，基于编程范式进行编程，可以快速搭建算子实现的代码框架。

Ascend C 编程范式把算子核内的处理程序，分成多个流水任务，以张量为数据载体，任务之间通过队列（Queue）进行通信和同步，并通过统一的内存管理模块（Pipe）管理任务间通信内存。

流水任务

流水任务指的是单核处理程序中主程序调度的并行任务。在核函数内部，可以通过流水任务实现数据的并行处理，进一步提升性能。以下面的流水任务示意图为例，单核处理程序的功能被拆分成 3 个流水任务：Stage1、Stage2、Stage3，每个任务专注于完成单一功能；需要处理的数据被切分成 n 片，使用 Progress1~n 表示，每个任务需要依次完成 n 个数据切片的处理。Stage 间的箭头表达数据间的依赖关系，比如 Stage1 处理完 Progress1 之后，Stage2 才能对 Progress1 进行处理。
2023-05-31-14-31-29.png
若 n=3，即待处理的数据被切分成 3 片，则上图中的流水任务运行起来的示意图如下，从运行图中可以看出，对于同一片数据，Stage1、Stage2、Stage3 之间的处理具有依赖关系，需要串行处理；不同的数据切片，同一时间点，可以有多个任务在并行处理，由此达到任务并行、提升性能的目的。

2023-05-31-14-32-21.png

矢量算子编程范式把算子的实现流程分为 3 个基本任务：CopyIn，Compute，CopyOut。CopyIn 负责搬入操作，Compute 负责矢量指令计算操作，CopyOut 负责搬出操作。
矩阵算子编程范式把算子的实现流程分为 5 个基本任务：CopyIn，Split，Compute，Aggregate，CopyOut。CopyIn 负责搬入操作，Split 负责数据切分操作，Compute 负责矩阵指令计算操作，Aggregate 负责数据汇聚操作，CopyOut 负责搬出操作。

矢量编程矩阵编程
2023-05-31-14-35-56.png2023-05-31-14-36-06.png

任务间通信和同步

不同的流水任务之间存在数据依赖，需要进行数据传递。
Ascend C 中使用 Queue 队列完成任务之间的数据通信和同步，提供 EnQue、DeQue 等基础 API。
Queue 队列管理不同层级的物理内存时，用一种抽象的逻辑位置（QuePosition）来表达各级别的存储，代替了片上物理存储的概念，达到隐藏芯片架构的目的。Queue 类型包括：VECIN、VECOUT、A1、A2、B1、B2、CO1、CO2，其中 VECIN、VECOUT 主要用于矢量编程，A1、A2、B1、B2、CO1、CO2 用于矩阵编程。

矢量编程

矢量编程中使用到的逻辑位置（QuePosition）定义如下：

  • 搬入数据的存放位置：VECIN；
  • 搬出数据的存放位置：VECOUT。

矢量编程主要分为 CopyIn、Compute、CopyOut 三个任务：

  • CopyIn 任务中将输入数据从 Global 内存搬运至 Local 内存后，需要使用 EnQue 将 LocalTensor 放入 VECIN 的 Queue 中；
  • Compute 任务等待 VECIN 的 Queue 中 LocalTensor 出队之后才可以完成矢量计算，计算完成后使用 EnQue 将计算结果 LocalTensor 放入到 VECOUT 的 Queue 中；
  • CopyOut 任务等待 VECOUT 的 Queue 中 LocalTensor 出队，再将其拷贝到 Global 内存。

2023-05-31-14-40-56.png

  1. Stage1：CopyIn 任务。
    • 使用 DataCopy 接口将 GlobalTensor 数据拷贝到 LocalTensor。
    • 使用 EnQue 将 LocalTensor 放入 VECIN 的 Queue 中。
  2. Stage2：Compute 任务。
    • 使用 DeQue 从 VECIN 中取出 LocalTensor。
    • 使用 Ascend C 接口完成矢量计算。
    • 使用 EnQue 将计算结果 LocalTensor 放入到 VECOUT 的 Queue 中。
  3. Stage3：CopyOut 任务。
    • 使用 DeQue 接口从 VECOUT 的 Queue 中去除 LocalTensor。
    • 使用 DataCopy 接口将 LocalTensor 拷贝到 GlobalTensor 上。

内存管理

任务间数据传递使用到的内存统一由内存管理模块 Pipe 进行管理。
如下图所示，Pipe 作为片上内存管理者，通过 InitBuffer 接口对外提供 Queue 内存初始化功能，开发者可以通过该接口为指定的 Queue 分配内存。
Queue 队列内存初始化完成后，需要使用内存时，通过调用 AllocTensor 来为 LocalTensor 分配内存，当创建的 LocalTensor 完成相关计算无需再使用时，再调用 FreeTensor 来回收 LocalTensor 的内存。

2023-05-31-14-50-04.png

1
2
3
4
5
6
7
8
9
10

TPipe pipe;
TQue<TPosition::VECOUT, 2> que;
uint8_t num = 4;
uint32_t len = 1024;
// 为TQue分配内存，分配内存块数为4，每块大小为1024Bytes
pipe.InitBuffer(que, num, len);
// AllocTensor分配Tensor长度为1024Bytes
LocalTensor<half> tensor1 = que.AllocTensor();
// 使用FreeTensor释放通过AllocTensor分配的Tensor
que.FreeTensor<half>(tensor1);

编程过程中使用到的临时变量内存同样通过 Pipe 进行管理。临时变量可以使用 TBuf 数据结构来申请指定 QuePosition 上的存储空间。使用 TBuf 申请的内存空间只能参与计算，无法执行 Queue 队列的入队出队操作。

1
2
3
4
5
6
7
8
9
10

// 为TBuf初始化分配内存，分配内存长度为1024字节
TPipe pipe;
TBuf<TPosition::VECCALC> calcBuf;
uint32_t len = 1024;
pipe.InitBuffer(calcBuf, len);

// 从calcBuf获取Tensor,Tensor大小为pipe分配的所有内存大小，即1024字节
LocalTensor<int32_t> tempTensor1 = calcBuf.Get<int32_t>();
// 从calcBuf获取Tensor,Tensor大小为128个int32_t类型元素的大小，即512字节
LocalTensor<int32_t> tempTensor1 = calcBuf.Get<int32_t>(128);

算子开发流程

Ascend C 算子开发分为快速流程和标准流程，快速开发流程为

  1. 完成算子核函数的开发；
  2. 基于内核调用符方式进行算子运行验证。

标注流程为：

  1. 完成算子核函数的开发；
  2. 完成单算子网络应用程序的开发；
  3. 基于 ACL 单算子调用方式进行算子运行验证。

基于 Ascend C 方式实现矢量算子的流程如下图所示：

2023-05-31-15-02-10.png

  • 算子分析：分析算子的数学表达式、输入、输出以及计算逻辑的实现，明确需要调用的 Ascend C 接口。
  • 核函数定义：定义 Ascend C 算子入口函数。
  • 根据矢量编程范式实现算子类：完成核函数的内部实现。

下面以张量相加为例：

算子分析

在开发算子代码之前需要分析算子的数学表达式、输入、输出以及计算逻辑的实现。

  1. 明确算子的数学表达式及计算逻辑。
    数学表达式为：z = x + y，其中 x 和 y 为 (8, 2048) 的张量，数据类型为 half，输入排列方式为 ND。
    计算逻辑：Ascend C 提供的矢量计算接口的操作元素都为 LocalTensor，输入数据需要先搬运进片上存储，然后使用计算接口完成两个输入参数相加，得到最终结果，再搬出到外部存储上。
  2. 明确输入和输出。
    • Add 算子有两个输入：x 与 y，输出为 z。
    • 算子输入的数据类型为 half（float16），算子输出的数据类型与输入数据类型相同。
    • 算子输入支持形状为（8，2048），输出形状与输入形状相同。
    • 算子输入支持的数据排列为：ND。
  3. 确定核函数名称和参数。
    • 自定义核函数名称，核函数命名为 add_tik2。
    • 根据对算子输入输出的分析，确定核函数有 3 个参数 x，y，z；
    • x，y 为输入在 Global Memory 上的内存地址，z 为输出在 Global Memory 上的内存地址。
  4. 确定算子实现所需接口。
    • 涉及外部存储和内部存储间的数据搬运，需要使用 DataCopy 来实；
    • 矢量计算的加法操作，使用双目指令 Add 接口 Add 实现 x+y；
    • 计算中使用到的 Tensor 数据结构为 LocalTensor，使用 Queue 队列进行管理，会使用到 EnQue、DeQue 等接口。

核函数定义

  1. 函数原型定义
    函数名为 add_tik2；根据算子分析中对算子输入输出的分析，确定有 3 个参数 x，y，z，参数类型统一设置成 uint8_t*，其中 x，y 都为输入内存，z 为输出内存；根据编写核函数核函数的规则介绍，返回值为 void，并增加 extern "C" 标识。除了需要按照 C/C++ 函数声明的方式定义核函数之外，还要为核函数加上额外的函数类型限定符和变量类型限定符。
    由此，可以得到函数原型定义为：

    1
2
3

    extern "C" __global__ __aicore__ void add_tik2(__gm__ uint8_t*x, __gm__ uint8_t* y, __gm__ uint8_t* z)
{
}

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

    1
2
3
4
5
6

    extern "C" __global__ __aicore__ void add_tik2(__gm__ uint8_t*x, __gm__ uint8_t* y, __gm__ uint8_t* z)
{
    KernelAdd op;
    op.Init(x, y, z);
    op.Process();
}

  3. 对核函数的调用进行封装，得到 add_tik2_do 函数，便于主程序调用。
    #ifndef __CCE_KT_TEST__表示该封装函数仅在编译运行 NPU 侧的算子时会用到，编译运行 CPU 侧的算子时，可以直接调用 add_tik2 函数。
    根据调用核函数章节，调用核函数时，除了需要传入参数 x，y，z，还需要传入 blockDim（核函数执行的核数）, l2ctrl（保留参数，设置为 nullptr）, stream（应用程序中维护异步操作执行顺序的 stream）来规定核函数的执行配置。

    1
2
3
4
5
6
7
8
9

    #ifndef __CCE_KT_TEST__

// call of kernel function
void add_tik2_do(uint32_t blockDim, void*l2ctrl, void* stream, uint8_t*x, uint8_t* y, uint8_t* z)
{
add_tik2<<<blockDim, l2ctrl, stream>>>(x, y, z);
}

# endif

算子类实现

2023-05-31-15-21-37.png

  1. 流水任务

    Add 算子的实现流程分为 3 个基本任务：CopyIn，Compute，CopyOut。

    • CopyIn 任务负责将 Global Memory 上的输入 Tensor xGm 和 yGm 搬运至 Local Memory，分别存储在 xLocal, yLocal；
    • Compute 任务负责对 xLocal, yLocal 执行加法操作，计算结果存储在 zLocal 中；
    • CopyOut 任务负责将输出数据从 zLocal 搬运至 Global Memory 上的输出 Tensor zGm 中。

  2. 数据通信

    • CopyIn，Compute 任务间通过 VECIN 队列 inQueueX，inQueueY 进行通信和同步，
    • Compute，CopyOut 任务间通过 VECOUT 队列 outQueueZ 进行通信和同步。

  3. 内存管理
    任务间交互使用到的内存、临时变量使用到的内存统一使用 pipe 内存管理对象进行管理。

算子类中主要实现上述流程，包含对外开放的初始化 Init 函数和核心处理函数 Process，Process 函数中会对上图中的三个基本任务进行调用；同时包括一些算子实现中会用到的私有成员，比如上图中的 Global Tensor 和 VECIN、VECOUT 队列等。KernelAdd 算子类具体成员如下：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22

class KernelAdd {
public:
    __aicore__ inline KernelAdd() {}
    // 初始化函数，完成内存初始化相关操作
    __aicore__ inline void Init(__gm__ uint8_t*x, __gm__ uint8_t* y, __gm__ uint8_t* 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 函数实现

    1. 获取该核函数需要处理的输入输出在 Global Memory 上的内存偏移地址。
      对于多核并行计算，需要把数据进行分片，分配到多个核上进行处理。Ascend C 核函数是在一个核上的处理函数，所以只处理部分数据，需要在初始化函数中获取该核函数需要处理的输入输出在 Global Memory 上的内存偏移地址。
      例如：数据整体长度 TOTAL_LENGTH 为 (8, 2048)，平均分配到 8 个核上运行，每个核上处理的数据大小 BLOCK_LENGTH 为 2048。block_idx 为核的逻辑 ID，(__gm__ half*)x + block_idx *BLOCK_LENGTH 即为单核处理程序中 x 在 Global Memory 上的内存偏移地址。
      对于单核上的处理数据，可以进行数据切块（Tiling），将数据切分成 8 块（并不意味着 8 块就是性能最优）。切分后的每个数据块再次切分成 2 块，即可开启 double buffer，实现流水线之间的并行。这样单核上的数据（2048 个数）被切分成 16 块，每块 TILE_LENGTH（128）个数据。上文代码表示 Pipe 为 inQueueX 分配了两块大小为 TILE_LENGTH * sizeof (half) 个字节的内存块，每个内存块能容纳 TILE_LENGTH（128）个 half 类型数据。
    2. 通过 Pipe 内存管理对象为输入输出 Queue 分配内存。
    1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18

    constexpr int32_t TOTAL_LENGTH = 8 * 2048;                            // total length of data
constexpr int32_t USE_CORE_NUM = 8;                                   // num of core used
constexpr int32_t BLOCK_LENGTH = TOTAL_LENGTH / USE_CORE_NUM;         // length computed of each core
constexpr int32_t TILE_NUM = 8;                                       // split data into 8 tiles for each core
constexpr int32_t BUFFER_NUM = 2;                                     // tensor num for each queue
constexpr int32_t TILE_LENGTH = BLOCK_LENGTH / TILE_NUM / BUFFER_NUM; // each tile length is seperated to 2 part, due to double buffer

__aicore__ inline void Init(__gm__ uint8_t*x, __gm__ uint8_t* y, __gm__ uint8_t*z)
{
    // get start index for current core, core parallel
    xGm.SetGlobalBuffer((__gm__ half*)x + block_idx *BLOCK_LENGTH);
    yGm.SetGlobalBuffer((__gm__ half*)y + block_idx *BLOCK_LENGTH);
    zGm.SetGlobalBuffer((__gm__ half*)z + block_idx *BLOCK_LENGTH);
    // pipe alloc memory to queue, the unit is 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));
}

  • Process 函数实现
    基于矢量编程范式，将核函数的实现分为 3 个基本任务：CopyIn，Compute，CopyOut。

    1
2
3
4
5
6
7
8
9
10
11

    __aicore__ inline void Process()
{
    // loop count need to be doubled, due to double buffer
    constexpr int32_t loopCount = TILE_NUM * BUFFER_NUM;
    // tiling strategy, pipeline parallel
    for (int32_t i = 0; i < loopCount; i++) {
        CopyIn(i);
        Compute(i);
        CopyOut(i);
    }
}

    • CopyIn 阶段

      1. 使用 DataCopy 接口将 GlobalTensor 数据拷贝到 LocalTensor。
      2. 使用 EnQue 将 LocalTensor 放入 VecIn 的 Queue 中。
      1
2
3
4
5
6
7
8
9
10
11
12

      __aicore__ inline void CopyIn(int32_t progress)
{
    // alloc tensor from queue memory
    LocalTensor<half> xLocal = inQueueX.AllocTensor<half>();
    LocalTensor<half> yLocal = inQueueY.AllocTensor<half>();
    // copy progress_th tile from global tensor to local tensor
    DataCopy(xLocal, xGm[progress * TILE_LENGTH], TILE_LENGTH);
    DataCopy(yLocal, yGm[progress * TILE_LENGTH], TILE_LENGTH);
    // enque input tensors to VECIN queue
    inQueueX.EnQue(xLocal);
    inQueueY.EnQue(yLocal);
}

    • Compute 阶段

      1. 使用 DeQue 从 VecIn 中取出 LocalTensor。
      2. 使用Ascend C 接口 Add 完成矢量计算。
      3. 使用 EnQue 将计算结果 LocalTensor 放入到 VecOut 的 Queue 中。
      4. 使用 FreeTensor 将释放不再使用的 LocalTensor。
      1
2
3
4
5
6
7
8
9
10
11
12
13
14

      __aicore__ inline void Compute(int32_t progress)
{
    // deque input tensors from VECIN queue
    LocalTensor<half> xLocal = inQueueX.DeQue<half>();
    LocalTensor<half> yLocal = inQueueY.DeQue<half>();
    LocalTensor<half> zLocal = outQueueZ.AllocTensor<half>();
    // call Add instr for computation
    Add(zLocal, xLocal, yLocal, TILE_LENGTH);
    // enque the output tensor to VECOUT queue
    outQueueZ.EnQue<half>(zLocal);
    // free input tensors for reuse
    inQueueX.FreeTensor(xLocal);
    inQueueY.FreeTensor(yLocal);
}

    • CopyOut 阶段

      1. 使用 DeQue 接口从 VecOut 的 Queue 中取出 LocalTensor。
      2. 使用 DataCopy 接口将 LocalTensor 拷贝到 GlobalTensor 上。
      3. 使用 FreeTensor 将不再使用的 LocalTensor 进行回收。
      1
2
3
4
5
6
7
8
9

      __aicore__ inline void CopyOut(int32_t progress)
{
    // deque output tensor from VECOUT queue
    LocalTensor<half> zLocal = outQueueZ.DeQue<half>();
    // copy progress_th tile from local tensor to global tensor
    DataCopy(zGm[progress * TILE_LENGTH], zLocal, TILE_LENGTH);
    // free output tensor for reuse
    outQueueZ.FreeTensor(zLocal);
}

  • double buffer
    double buffer 通过将数据搬运与矢量计算并行执行以隐藏数据搬运时间并降低矢量指令的等待时间，最终提高矢量计算单元的利用效率。

    为减少 Vector 等待时间，double buffer 机制将待处理的数据一分为二，比如 Tensor1、Tensor2。

    1. 当 Vector 对 Tensor1 中数据进行 Compute 时，Tensor2 可以执行 CopyIn 的过程；
    2. 而当 Vector 切换到计算 Tensor2 时，Tensor1 可以执行 CopyOut 的过程。

本帖最后由 匿名用户 于 2023/08/23 20:07:40 编辑

我要发帖子

我们使用cookie来确保您的高速浏览体验。继续浏览本站，即表示您同意我们使用cookie。查看详情