Ascend C的编程模型

news2024/11/24 20:12:24

1 并发执行

Ascend C和cudnn相似，都是一种多核心编程的范式。想要了解Ascend C，必须得先掌握这种“多核”是怎么实现得。

多核执行，说白了就是使用CPU/GPU/Ascend的物理多核并发去执行一段流程，一般情况下，可以通过以下几种方式实现：

多线程并行处理：使用多线程可以将循环中的任务分配给多个线程同时执行，提高代码的执行效率。可以使用线程池来管理线程的创建和销毁，避免频繁创建和销毁线程的开销。
多进程并行处理：使用多进程可以将循环中的任务分配给多个进程同时执行，充分利用多核处理器的优势。可以使用multiprocessing模块来创建和管理多个进程，并通过进程间通信来实现数据的共享和同步。
向量化操作：对于一些数值计算密集型的任务，可以使用向量化操作来提高代码的执行效率。向量化操作利用了现代处理器的SIMD（Single Instruction, Multiple Data）指令集，可以同时对多个数据进行相同的操作，减少了循环的开销。在Python中，可以使用NumPy库来进行向量化操作。
并行计算框架：使用并行计算框架可以将循环中的任务分布到多个计算节点上并行执行，提高代码的执行速度。常见的并行计算框架包括Apache Hadoop、Apache Spark等。这些框架提供了分布式计算的能力，可以将任务分布到多台计算机上进行并行计算。
编译器指令OpenMP：是一种支持多平台共享内存并行编程的API，它提供了一组编译器指令和库函数，使得开发人员能够方便地将现有代码并行化，以利用多个核心或处理器进行并行执行。OpenMP支持多种编程语言和操作系统，具有易用性、可移植性和灵活性等特点。通过使用OpenMP的并行化指令，如#pragma omp parallel和#pragma omp for，可以轻松地将代码块或循环并行化，并可以通过设置线程数量来控制并行执行的程度。OpenMP还提供了特定的编译器指令来处理并发任务和循环并行化。例如，#pragma omp parallel用于创建一组线程来并行执行指定的代码块，而#pragma omp for用于并行化循环。这些指令允许开发人员精细控制并行化的程度，包括设置线程数量和使用特定的子句来指定条件并行、数据处理等。

2 Ascend C编程模型

Ascend C算子编程是SPMD（Single-Program Multiple-Data）编程。假设，从输入数据到输出数据需要经过3个阶段任务的处理（T1、T2、T3）。如下图所示，SPMD会启动一组进程，并行处理待处理的数据。对待处理数据切分，把切分后数据分片分发给不同进程处理，每个进程对自己的数据分片进行3个任务的处理。
SPMD模型
具体到Ascend C编程模型中的应用，是将需要处理的数据被拆分并同时在多个计算核心（类比于上文介绍中的多个进程）上运行，从而获取更高的性能。多个AI Core共享相同的指令代码，每个核上的运行实例唯一的区别是block_idx不同，每个核通过不同的block_idx来识别自己的身份。block的概念类似于上文中进程的概念，block_idx就是标识进程唯一性的进程ID。并行计算过程的示意图如下图所示。
AiCore并行计算示意图
下面的代码片段取自于Ascend CAdd算子的实现代码，算子被调用时，所有的计算核心都执行相同的实现代码，入口函数的入参也是相同的。每个核上处理的数据地址需要在起始地址上增加GetBlockIdx()*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
        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 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));
    }
    ...
}

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