__gloabl__ __aircore__ void kernel_name(argument list);

#define GM_ADDR __gm__ unit8_t* __restrict__

function_name(argument list);

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

HelloWorld<<<8, nullptr, stream>>>(fooDevice));

aclError aclrtSynchronizeStream(aclrtStream stream);

Ascend C的并行编程式范式核心要素

• 一组并行计算任务
• 通过队列实现任务之间的通信和同步
• 程序员自主表达对并行计算任务和资源的调度

典型的计算范式

• 基本的矢量编程范式：计算任务分为CopyIn,Compute,CopyOut
• 基本的矩阵编程范式：计算任务分为CopyIn,Compute,Aggregate,CopyOut
• 复杂的矢量/矩阵编程范式，通过将矢量/矩阵的Out/ln组合在一起的方式来实现复杂计算数据流

流水任务

流水任务（Stage）指的是单核处理程序中主程序调度的并行任务。

在核函数内部，可以通过流水任务实现数据的并行处理来提升性能

举例来说，单核处理程序的功能可以拆分为3个流水任务：Stage1、Stage2、Stage3，每个任务专注数据切片的处理。Stage间的剪头表达数据间的依赖，比如Stage1处理完Progress1之后，Stage2才能对Proress1进行处理。

若Progres的n=3，待处理的数据被切分成3片，对于同一片数据，Stage1、Stage2、Stage3之间的处理具有依赖关系，需要串行处理；不同的数据切片，同一时间点，可以有多个流水任务Stage在并行处理，由此达到任务并行、提升性能的目的

任务间通信和同步

数据通信与同步的管理者

Ascend C中使用Queue队列完成任务之间的数据通信和同步，Queue提供了EnQue、DeQue等基础API。

Queue队列管理NPU上不同层级的物理内存时,用一种抽象的逻辑位置(QuePosition)来表达各个级别的存储(Storage Scope),代替了片上物理存储的概念,开发者无需感知硬件架构。

矢量编程中Queue类型(逻辑位置)包括:VECIN、VECOUT

数据的载体

Ascend C使用GlobalTensor和LocalTensor作为数据的基本操作单元，它是各种指令API直接调用的对象,也是数据的载体

矢量编程任务间通信和任务

矢量编程中的逻辑位置(QuePosition)：搬入数据的存放位置：VECIN、搬出数据的存放位置：VECOUT。

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

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

Stage1：CopyIn任务

使用DataCopy接口将GlobalTensor拷贝纸LocalTensor

使用EnQue将LocalTensor放入VECIN的Queue中

Stage2：Compute任务

使用DeQue从VECIN中取出LocalTensor

使用Ascend C指令API完成矢量计算:Add

使用EnQue将结果LocalTensor放入VECOUT的Queue中

Stage3：CopyOut任务

使用DeQue接口从VECOUT的Queue中取出LocalTensor

使用DataCopy接口将LocalTensor拷贝至GlobalTensor

内存管理

任务见数据传递使用到的内存统一由内存管理模块Pipe进行管理。

Pipe作为片上内存管理者，通过InitBuffer接口对外提供Queue内存初始化功能，开发者可以通过该接口为指定的Queue分配内存。

Queue队列内存初始化完成后，需要使用内存时，通过调用AllocTensor来为LocalTensor分配内存给Tensor，当创建的LocalTensor完成相关计算无需再使用时，再调用FreeTensor来回收LocalTensor的内存

临时变量内存管理

编程过程中使用到的临时变量内存同样通过Pipe进行管理。临时变量可以使用TBuf数据结构来申请指定QuePosition上的存储空间，并使用Get()来将分配到的存储空间分配给新的LocalTensor从TBuf上获取全部长度，或者获取指定长度的LocalTensor

LocalTensor<T> Get<T>();
LocalTensor<T> Get<T>(uint32_t len);

Tbuf及Get接口的示例

//为TBuf初始化分配内存，分配内存长度为1024字节
TPipe pipe;
TBuf<TPosition::VECIN> calcBuf; //模板参数为QuePosition中的VECIN类型
uint32_t byteLen = 1024;
pipe.InitBuffer(calcBuf,byteLen);
//从calcBuf获取Tensor,Tensor为pipe分配的所有内存大小，为1024字节
LocalTensor<int32_t> tempTensor1 = calcBuf.Get<int32_t>();
//从calcBuf获取Tensor,Tensor为128个int32_t类型元素的内存大小，为512字节
LocalTensro<int32_t> tempTensor1 = calcBuf.Get<int32_t>(128);

使用TBuf申请的内存空间只能参与计算，无法执行Queue队列的入队出队操作

Ascend C矢量编程

算子分析

开发流程

算子分析：分析算子的数学表达式、输入、输出以及计算逻辑的实现，明确需要调用的Ascend接口

核函数定义：定义Ascend算子入口函数

根据矢量编程范式实现算子类：完成核函数的内部实现

以ElemWise(ADD)算子为，数学公式

为简单起见，设定张量x,y,z为固定shape(8,2048),数据类型dtype为half类型，数据排布类型format为ND,核函数名称为add_custom

 

算子分析

明确算子的数学表达式及计算逻辑

Add算子的数学表达式为

 

明确输入输出

Add算子有两个：

输入数据类型为half，输出数据类型与输入数据类型相同。输入支持固定shape(8,2048),输出shape与输入shape相同，输入数据排布类型为ND

 

确定核函数名称和参数

自定义核函数明，如add_custom,根据输入输出，确定核函数有3个入参x,y,z
x,y为输入在GlobalMemory上的内存地址，z为输出在globalMemory上的内存地址

确定算子实现所需接口

涉及内外部存储间的数据搬运，使用数据搬移接口：DataCopy实现

涉及矢量计算的加法操作，使用矢量双目指令：Add实现

使用到LocalTensor,使用Queue队列管理，会使用到Enque，Deque等接口。

算子实现

核函数定义

在add_custom核函数的实现中示例化KernelAdd算子类，调用Init()函数完成内存初始化，调用Process()函数完成核心逻辑。

注：算子类和成员函数名无特殊要求，开发者可根据自身的C/C++编码习惯，决定核函数中的具体实现。

// implementation of kenel function
extern "C" __global__ __aicore__ void add_custom(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z)
{
	kernelAdd op;
	op.Init(x,y,z);
	op.Process();
}

对于核函数的调用，使用内置宏__CCE_KT_TEST__来标识<<<…>>>仅在NPU模式下才会编译到(CPU模式g++没有<<<…>>>的表达),对核函数的调用进行封装，可以在封装函数中补充其他逻辑，这里仅展示对于核函数的调用。

#ifndef __CCE_KT_TEST__
// call of kernel function
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);
}

算子类实现

CopyIn任务：将Global Memory上的输入Tensor xGm和yGm搬运至Local Memory,分别存储在xlocal，ylocal。

Compute任务：对xLocal，yLocal执行加法操作，计算结果存储在zlocal中。

CopyOut任务：将输出数据从zlocal搬运至Global Memory上的输出tensor zGm中。

CopyIn.Compute任务间通过VECIN队列和inQueueX,inQueueY进行通信和同步。

Compute，CopyOut任务间通过VECOUT和outQueueZ进行通信和同步。

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

向量加法z=x+y 代码样例 TPIPE流水式编程范式

算子类实现

算子类类名: KernelAdd

初始化函数Init()和核心处理函数Process()

三个流水任务：CopyIn(),Compute(),CopyOut()

Process的含义

TQue模板的BUFFER)NUM的含义：

该Queue的深度，double buffer优化技巧

class KernelAdd{
public:
	__aicore__ inline KernelAdd()
	//初始化函数，完成内存初始化相关操作
	__aicore__ inline voide 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 process){}
	// 计算函数，完成Compute阶段的处理，被Process函数调用
	__aicore__ inline void Compute(int32_t process){}
	// 搬出函数，完成CopyOut阶段的处理，被Process函数调用
	__aicore__ inline void CopyOut(int32_t process){}

private:
	// pipe内存管理对象
	TPipe pipe;
	// 输入数据Queue队列管理对象，QuePosition为VECIN
	TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX, inQueueY;
	// 输出数据Queue队列管理对象，QuePosition为VECOUT
	TQue<QuePosition::VECOUT， BUFFER_NUM> outQueueZ;
	// 管理输入输出的Global Memory内存地址的对象，其中xGm,yGm为输入,zGm为输出
	GlobalTensor<half> xGm, yGm ,zGm;
};

Init()函数实现

使用多核并行计算，需要将数据切片，获取到每个核实际需要处理的在Global Memory上的内存偏移地址。

数据整体长度TOTAL_LENGTH为8 * 2048，平均分配到8个核上运行，每个核上处理的数据大小BLOCK_LENGTH为2048，block_idx为核的逻辑ID，(gm half*)x + GetBlockIdx() *
BLOCK_LENGTH即索引为block_idx的核的输入数据在Global Memory上的内存偏移地址

对于单核处理数据，可以进行数据切块（Tiling）,将数据切分成8快，切分后的每个数据块再次切分成BUFFER_NUM=2块，可开启double buffer，实现流水线之间的并行。

单核需要处理的2048个数据切分成16块，每块TILE_LENGTH=128个数据，Pipe为inQueueX分配了BUFFER_NUM块大小为TITLE_LENGTH * sizeof(half)个字节的内存块，每个内存块能容纳TILE_LENGTH=128个half类型数据

代码示例

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 ccore
constexpr int32_t TILE_NUM = 8; //split data into 8 tiles
constexpr int32_t BUFFER_NUM = 2; //tensor num for each queue
constexpr int32_t TILE_LENGTH = BLOCK_LENGTH / TILE_NUM / BUFFER_NUM; //seperate to 2 parts, due to double buffer

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

Process()函数实现

代码示例

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

__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(xLocal,yGm[progress * TILE_LENGTH], TILE_LENGTH);
	// enque input tensors to VECIN queue
	inQueueX.EnQue(xLocal);
	inQueueY.EnQue(yLocal);
}

__aicore__ inline void Compute(int32_t progress)
{
	//dque 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)l
	// free input tensors for reuse
	inQueueX.FreeTensor(xLocal);
	inQueueY.FreeTensor(yLocal);
}

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

double buffer机制

double buffer通过将数据搬运与矢量计算并执行以隐藏数据搬运时间并降低矢量指令的等待时间，最终提高矢量计算单元的利用效率1个Tensor同一时间只能进行搬入、计算和搬出三个流水任务中的一个，其他两个流水任务涉及的硬件但愿则处于Idle状态。

如果将待处理的数据一分为而，比如Tensor1、Tensor2。

• 当矢量计算单元对于Tensor1进行Compute时，Tensor2可以进行CopyIn的任务
• 当矢量计算单元对于Tensor2进行Compute时，Tensor1可以进行CopyOut的任务
• 当矢量计算单元对于Tensor2进行CopyOut时，Tensor2可以进行CopyIn的任务
  由此，数据的进出搬运和矢量计算之间实现你并行，硬件单元闲置问题得以有效缓解

Ascend C 算子调用

HelloWorld样例

运行CPU模式包含的头文件

运行NPU模式包含的头文件

核函数的定义

内置宏__CE_KT_TEST__：区分运行CPU模式或NPU模式逻辑的标志

主机侧执行逻辑：负责数据在主机侧内存的申请，主机到设备的拷贝，核函数执行同步和回收资源的工作

设备侧执行逻辑

主机侧执行CPU模式逻辑：使用封装的执行宏ICPU_RUN_KF

主要包括:

gMAlloc(…)：申请CPU模式下的内存空间

ICPU_RUN_KF：使用封装的执行宏

GmFree：释放CPU模式下的内存空间

流程

AscendCL初始化—>运行管理资源申请—>Host数据传输至Device—>执行任务并等待—>Device数据传输至Host—>运行资源释放—>AscendCL去初始化

主机侧执行NPU模式逻辑：使用内核调用符<<<…>>>

重要接口

• aclInit
• aclCreateStream
• aclMallocHost
• aclMalloc
• aclMemcpy
• <<<…>>>
• aclrtSynchronizeStream
• aclrtFree
• aclrtfreeHost
• aclrtDestoryStream
• aclFinalize

AddCustom样例

Ascend C矢量算子样例代码

1. 核函数源文件：add_custom.app
2. 真值数据生成脚本：add_custom.py
3. CmakeLists.txt：方便对多个源文件进行编译
4. 读写数据文件辅助函数：data_utils.h
5. 主机侧源文件：main.cpp
6. 一键执行脚本：run.sh
7. 组织CPU模式和NPU模式下编译的cmake脚本

点击关注，第一时间了解华为云新鲜技术~