手敲，Ascend 算子开发入门笔记分享

本文分享自华为云社区《Ascend算子开发入门笔记》，作者： JeffDing 。

基础概念

什么是 Ascend C

Ascend C 是 CANN 针对算子开发场景推出的编程语言,原生支持 C 和 C++标准规范，最大化匹配用户开发习惯；通过多层接口抽象、自动并行计算、孪生调试等关键技术，极大提高算子开发效率，助力 AI 开发者低成本完成算子开发和模型调优部署。

使用 Ascend C 开发自定义算子的优势

• C/C++原语编程，最大化匹配用户的开发习惯

• 编程模型屏蔽硬件差异，编程范式提高开发效率

• 多层级 API 封装，从简单到灵活，兼顾易用与高效

• 孪生调试，CPU 侧模拟 NPU 侧的行为，可优化在 CPU 侧调试

昇腾计算架构 CANN

CANN 介绍网站：https://www.hiascend.com/software/cann

AI Core 是 NPU 卡的计算核心,NPU 内部有多个 AI Core。每个 AI Core 相当于多核 CPU 中的一个核心

SIMD

SIMD，也就是单指令多数据计算，一条指令可以处理多个数据：Ascend C 编程 API 主要是向量计算 API 和矩阵运算 API,计算 API 都是 SIMD 样式

并行计算之 SPMD 数据并行与流线型并行

SPMD 数据并行原理

• 启动一组进程，他们运行的相同程序

• 把待处理数据切分，把切分后数据分片分发给不同进程处理

• 每个进程对自己的数据分片进行 3 个任务 T1、T2、T3 的处理

流水线并行原理

• 启动一组进程

• 对数据进行切分

• 每个进程都处理所有的数据切片，对输入数据分片只做一个任务的处理

Ascend C 编程模型与范式

并行计算架构抽象

使用 Ascend C 编程语言开发的算子运行在 AI Core 上，AI Core 是昇腾 AI 处理器中的计算核心一个 AI 处理器内部有多个 AI Core，AI Core 中包含计算单元、存储单元、搬运单元等核心组件

计算单元包括了三种基础计算资源

1. Scalar 计算单元：执行地址计算、循环控制等标量计算工作，并把向量计算、矩阵计算、数据半圆、同步指令发射给对应单元执行

2. Cube 计算单元：负责执行矩阵运算

3. Vector 计算单元：负责执行向量计算

搬运单元负责在 Global Memory 和 Local Memory 之间搬运数据，包含搬运单元 MTE(Memory Transfer Engine,数据搬入单元),MTE3（数据搬出单元）

存储单元为 AI Core 的内部存储，统称为 Local Memory 与此相对应，AI Core 的外部存储称之为 Global Memory

异步指令流

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

同步信号流

指令间可能存在依赖关系，为了保证不同指令队列间的指令按照正确的逻辑关系执行，Scalar 计算单元也会给对应单元下发同步指令

计算数据流

DMA 搬入单元把数据搬运到 Local Memory，Vector/Cube 计算单元完成数据计算，并把计算结构写回 Local Memory，DMA 搬出单元把处理好的数据搬运回 Global Memory

SPMD 编程模型介绍

Ascend C 算子编程是 SPMD 的编程，将需要处理的数据拆分并行分布在多个计算核心上运行多个 AI Core 共享相同的指令代码，每个核上的运行实例唯一的区别是 block_idx 不同 block 的类似于进程，block_idx 就是标识进程唯一性的进程 ID,编程中使用函数 GetBlockIdx()获取 ID

核函数编写及调用

核函数（Kernel Function）是 Acend C 算子设备侧的入口。Ascend C 允许用户使用核函数这种 C/C++函数的语法扩展来管理设备侧的运行代码，用户在核函数中实现算子逻辑的编写，例如自定义算子类及其成员函数以实现该算子的所有功能。核函数是主机侧和设备侧连接的桥梁

核函数是直接在设备侧执行的代码。在核函数中，需要为在一个核上执行的代码规定要进行的数据访问和计算操作，SPMD 编程模型允许核函数调用时，多个核并行地执行同一个计算任务。

使用函数类型限定符

除了需要按照 C/C++函数声明的方式定义核函数之外，还要为核函数加上额外的函数类型限定符，包含__global__和__aicore__

使用__global__函数类型限定符来标识它是一个核函数，可以被<<<…>>>调用；使用__aicore__函数类型限定符来标识该函数在设备侧 AI Core 上执行

__gloabl__ __aircore__ void kernel_name(argument list);

使用变量类型限定符

为了方便：指针入参变量统一的类型定义为__gm__uint8_t*

用户可统一使用 uint8_t 类型的指针，并在使用时转化为实际的指针类型；亦可直接传入实际的指针类型

规则或建议

1. 核函数必须具有 void 返回类型

2. 仅支持入参为指针类型或 C/C++内置数据类型(Primitive Data Types),如：half* s0、flat* s1、int32_t c

3. 提供了一个封装的宏 GM_ADDR 来避免过长的函数入参列表

#define GM_ADDR __gm__ unit8_t* __restrict__

调用核函数

核函数的调用语句是 C/C++函数调用语句的一种扩展

常见的 C/C++函数调用方式是如下的形式：

function_name(argument list);

核函数使用内部调用符<<<…>>>这种语法形式，来规定核函数的执行配置：

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

注：内核调用符仅可在 NPU 模式下编译时调用,CPU 模式下编译无法识别该符号

blocakdim,规定了核函数将会在几个核上执行,每个执行该核函数的核会被分配一个逻辑 ID，表现为内置变量 block_idx,编号从 0 开始，可为不同的逻辑核定义不同的行为，可以在算子实现中使用 GetBlockIDX()函数来获得。

l2ctl,保留函数，展示设置为固定值 nullptr。

stream：类型为 aclrtStream,stream 是一个任务队列，应用程序通过 stream 来管理任务的并行

使用内核调用符<<<…>>>调用核函数：

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

blockDim 设置为 8，表示在 8 个核上调用了 HelloWorld 核函数，每个核都会独立且并行地执行该核函数 Stream 可以通过 aclrtCreateStream 来创建，它的作用是在当前进程或线程中显式创建一个 aclrtStream argument list 设置为 cooDevice 这 1 个入参。

核函数的调用是异步的，核函数的调用结束后，控制权立刻返回给主机侧。

强制主机侧程序等待所有核函数执行完毕的 API(阻塞应用程序运行，直到指定 Stream 中的所有任务都完成，同步接口）为 aclrtSynchronizeStream

aclError aclrtSynchronizeStream(aclrtStream stream);

编程 API 介绍

Ascend C 算子采用标准 C++语法和一组类库 API 进行编程。

计算类 API：标量计算 API、向量计算 API、矩阵计算 API、分别实现调用 Scalar 计算单元、Vector 计算单元、Cube 计算单元。

数据搬运 API：基于 Local Memory 数据进行计算、数据需要先从 Gloabl Memory 搬运至 Local Memory，再使用计算接口完成计算，最后从 Local Memory 搬出至 Gloabl Memory。比如 DataCopy 接口。

内存管理 API：用于分配管理内存，比如 AllocTensor、FreeTensor 接口。

任务同步 API：完成任务间的通信和同步，比如 EnQue、DeQue 接口。不同的指令异步并行执行，为了保证不同指令队列间的指令按照正确的逻辑关系执行，需要向不同的组件发送同步指令。

Ascend C API 用于计算的基本数据类型都是 Tensor：GlobalTensor 和 LocalTensor。

4 级 API 定义

4 级 API 定义：API 根据用户使用的场景分为 4 级。

3 级 API,运算符重载,支持+, - ,* ,/ ,= ,| ,& ,^ ,> ,< ,>- ,<= 实现计算的简单表述，类似 dst=src1+src2。

2 级连续计算 API，类似 Add(dst,src1,src2,count),针对源操作数的连续 COUNT 个数据进行计算连续写入目的操作数，解决一维 tensor 的连续 count 个数据的计算问题。

1 级 slice 计算 API，解决多维数据中的切片计算问题（开发中）。

0 级丰富功能计算 API，可以完整发挥硬件优势的计算 API，该功能可以充分发挥 CANN 系列芯片的强大指令,支持对每个操作数的 repeattimes,repetstride,MASK 的操作。调用类似：Add(dst,src1,src2,repeatTimes,repeatParams);

流水编程范式介绍

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

• 快速开发编程的固定步骤

• 统一代码框架的开发捷径

• 使用者总结出的开发经验

• 面向特定场景的编程思想

• 定制化的方法论开发体验

抽象编程模型“TPIPE 并行计算"

针对各代 Davinci 芯片的复杂数据流，根据实际计算需求，抽象出并行编程范式，简化流水并行。

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,zx,y 为输入在 GlobalMemory 上的内存地址，z 为输出在 globalMemory 上的内存地址

确定算子实现所需接口

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

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

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

算子实现

核函数定义

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

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

// implementation of kenel functionextern "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 functionvoid 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 dataconstexpr int32_t USE_CORE_NUM = 8;  //num of core usedconstexpr int32_t BLOCK_LENGTH = TOTAL_LENGTH / USE_CORE_NUM;  //length computed of each ccoreconstexpr int32_t TILE_NUM = 8; //split data into 8 tilesconstexpr int32_t BUFFER_NUM = 2; //tensor num for each queueconstexpr 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 脚本

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