文档
注册
评分
提单
论坛
小AI

SPMD模型

Ascend C算子编程是SPMD(Single-Program Multiple-Data)编程SPMD是一种常用的并行计算的方法,是提高计算速度的有效手段。

假设,从输入数据到输出数据需要经过3个阶段任务的处理(T1、T2、T3)。如下图所示,SPMD模式下,系统会启动一组进程,并行处理待处理的数据:首先待处理数据会被切分成多个数据分片,切分后的数据分片随后被分发给不同进程处理,每个进程接收到一个或多个数据分片,并独立地对这些分片进行3个任务的处理。

图1 SPMD数据并行示意图

具体到Ascend C编程模型中的应用,是将需要处理的数据被拆分并同时在多个计算核心(类比于上文介绍中的多个进程)上运行,从而获取更高的性能。多个AI Core共享相同的指令代码,每个核上的运行实例唯一的区别是block_idx不同,每个核通过不同的block_idx来识别自己的身份。block的概念类似于上文中进程的概念,block_idx就是标识进程唯一性的进程ID。并行计算过程如下图所示。

图2 SPMD并行计算示意图

下面的代码片段取自于Ascend C Add算子的实现代码,算子被调用时,所有的计算核心都执行相同的实现代码,入口函数的入参也是相同的。每个核上处理的数据地址需要在起始地址上增加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();
}
搜索结果
找到“0”个结果

当前产品无相关内容

未找到相关内容,请尝试其他搜索词