ascend c算子开发学习笔记

Murphy_232

已于 2024-07-24 13:13:33 修改

阅读量280

点赞数 4

文章标签： c++ 性能优化华为云

于 2024-07-24 11:27:06 首次发布

本文链接：https://blog.csdn.net/m0_69834888/article/details/140658386

版权

基础架构

Ascend C算子运行在AI Core上，包括三种单元：

1. 计算单元

Scalar（标量）计算单元

负责各类标量数据运算以及程序的流程控制，主要用于发射指令

Vector（向量）计算单元

执行向量计算，类似于SIMD指令，可以在一个周期里面完成多个数据的计算

Cube（矩阵）计算单元

负责执行矩阵运算，一条指令可以完成一个矩阵的运算

2. 存储单元

AI Core自己有内部的存储（Local Memory），与之对应的（例如DDR内存中）为外部存储（Global Memory），需要进行运算时，将数据从GM搬运到LM中。

3. 控制单元

提供指令控制，负责整个AI Core的运行。

硬件抽象架构

AI Core内部的异步并行计算过程：Scalar计算单元读取指令序列，并把向量计算、矩阵计算、数据搬运指令发射给对应单元的指令队列，向量计算单元、矩阵计算单元、数据搬运单元异步的并行执行接收到的指令。

SPMD模型

SPMD（Single-Program Multiple-Data）是一种并行计算的方法，将要处理的数据拆分给多个进程并行执行，每个进程都执行完整的处理操作

具体到Ascend C中，就是将数据切分给不同的AI Core来并行执行。

核函数

使用Ascend C进行编程时，我们编写一份算子实现代码，算子被调用时，将启动N个运行示例，在N个核上运行。

核函数（Kernel Function）是Ascend C算子设备侧实现的入口。在核函数中，需要为在一个核上执行的代码规定要进行的数据访问和计算操作，当核函数被调用时，多个核都执行相同的核函数代码，具有相同的参数，并行执行。

host侧与device侧

启动进程的一侧为host，执行任务的一侧是device

核函数定义调用关系

核函数定义：

extern "C" __global__ __aicore__ void add_custom(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z);

核函数调用：

kernel_name<<<blockDim, l2ctrl, stream>>>(argument list);

blockDim，规定了核函数将会在几个核上执行；

l2ctrl，保留参数，暂时设置为固定值nullptr；

stream，用于维护一些异步操作的执行顺序

简单示例

// 实现核函数
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();
}

// 调用核函数
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);
}

编程范式

AICore内部的执行单元是一个异步并行的过程，每一个执行单元都可以看成是流水线上的节点，通过流水并行的方式来提高计算效率。

如上图所示，Vector编程范式把算子的实现流程分为3个基本任务：CopyIn，Compute，CopyOut。

CopyIn搬入：将输入数据从GM搬运到LocalMemory（VECIN用于表达矢量计算搬入数据的存放位置），完成搬运后执行入队列操作；
Compute向量计算：完成队列出队后，从LocalMemory获取数据并计算，计算完成后执行入队操作；
CopyOut搬出：完成队列出队后，将计算结果从LocalMemory（VECOUT用于表达矢量计算搬出数据的存放位置）搬运到GM。

具体流程如下：

Pipe pipe;   //创建全局的资源管理   
TQue<VecIn, 1> queIn;  //创建CopyIn阶段的队列
TQue<VecOut, 1> queOut; //创建CopyOut阶段的队列
// Init 阶段：
pipe.InitBuffer(queIn, 2, 1024);  // 开启doublebuffer
for-loop {
    //CopyIn 阶段{
    auto tensor = queIn.AllocTensor<half>();     //从Que上申请资源, 长度1024
    DataCopy(tensor, gm, len);                   //搬运数据从GM到VECIN
    queIn.EnQue(tensor); 
    }
    //Compute阶段{
    auto tensor = queIn.DeQue<half>();
    auto tensorOut = queOut.AllocTensor<half>();
    Abs(tensorOut, tensor, 1024);
    queIn.FreeTensor(tensor);
    queOut.EnQue(tensorOut);
    }
    //CopyOut 阶段{
    auto tensor = queOut.DeQue<half>();
    DataCopy(gmOut, tensor, 1024);
    queOut.FreeTensor(tensor);
    }
}

任务间数据传递使用到的内存、事件等资源统一由管理模块Pipe进行管理。如下所示的内存管理示意图，TPipe通过InitBuffer接口对外提供Queue内存初始化功能，开发者可以通过该接口为指定的Queue分配内存。