2023年5月6日，在昇腾AI开发者峰会上，华为正式发布了面向算子开发场景的昇腾AscendC编程语言。AscendC原生支持C/C++编程规范，通过多层接口抽象、并行编程范式、孪生调试等技术，极大提高了算子的开发效率，帮助AI开发者低成本完成算子开发和模型调优部署。
1昇腾AI软硬件基础
和CUDA开发的算子运行在GPU上一样，基于AscendC开发的算子，可以通过异构计算架构CANN运行在昇腾AI处理器上。CANN是使能昇腾AI处理器的一个软件栈，通过软硬件协同优化，能够充分发挥昇腾AI处理器的强大算力。从下面的架构图可以清楚的看到，使用AscendC编程语言开发的算子通过编译器编译和运行时调度，最终运行在昇腾AI处理器上。
我们知道，通用计算就是我们常写的一些在CPU上运行的计算，它擅长逻辑控制和串行计算，而AI计算相对通用计算来说，更擅长并行计算，可支持大规模的计算密集型任务。如下面左图所示，做一个矩阵乘，使用CPU计算需要三层for循环，而右图在昇腾AI处理器上使用vector计算单元，只需要两层for循环，最小计算代码能同时计算多个数据的乘加，更近一步，如果使用Cube计算单元，只需要一条语句就能完成一个矩阵乘的计算，这就是我们所说的SIMD。因此，我们通常使用AI处理器来进行大量的并行计算。
NPU不能独立运行，需要与CPU协同工作，可以看成是CPU的协处理器，CPU负责整个操作系统运行，管理各类资源并进行复杂的逻辑控制，而NPU主要负责并行计算任务。在基于CPU+NPU的异构计算架构中，NPU与CPU通过PCIe总线连接在一起来协同工作，CPU所在位置称为主机端，而NPU所在位置称为设备端，示意图如下：
这里再详细介绍一下昇腾AI处理器。昇腾AI处理器有不同的型号和产品形态，小到模块、加速卡，大到服务器、集群。昇腾AI处理器里面最核心的部件是AICore，有多个，是神经网络加速的计算核心，每一个AICore就相当于我们大家平时理解的多核cpu里的每个核，使用AscendC编程语言开发的算子就运行在AICore上，因为核心的神经网络计算的加速都来源于AICore的算力。
AICore内部的并行计算架构抽象如下图所示：
这个并行计算架构抽象核心包含了几个大的部件，AICore外面有一个GobalMemory，是多个AICore共享的，在AICore内部有一块本地内存LocalMemory，因为靠近计算单元，所以它的带宽会非常高，相对的容量就会很小，比如一般是几百K到1M。AICore内部的核心组件有三个计算单元，标量计算单元、向量计算单元，矩阵计算单元。另外还有一个DMA搬运单元，DMA搬运单元负责在GlobalMemory和LocalMemory之间搬运数据。
AICore内部的异步并行计算过程：Scalar计算单元读取指令序列，并把向量计算、矩阵计算、数据搬运指令发射给对应单元的指令队列，向量计算单元、矩阵计算单元、数据搬运单元异步并行执行接收到的指令。该过程可以参考上图中蓝色箭头所示的指令流。不同的指令间有可能存在依赖关系，为了保证不同指令队列间的指令按照正确的逻辑关系执行，Scalar计算单元也会给对应单元下发同步指令。各单元之间的同步过程可以参考上图中的橙色箭头所示的同步信号流。
AICore内部数据处理的基本过程：DMA搬入单元把数据搬运到LocalMemory，Vector/Cube计算单元完成数据，并把计算结果写回LocalMemory，DMA搬出单元把处理好的数据搬运回GlobalMemory。该过程可以参考上图中的红色箭头所示的数据流。
2AscendC编程模型基础
2.1AscendC编程范式
AscendC编程范式是一种流水线式的编程范式，把算子核内的处理程序，分成多个流水任务，通过队列完成任务间通信和同步，并通过统一的内存管理模块管理任务间通信内存。流水编程范式应用了流水线并行计算方法。
若n=3，即待处理的数据被切分成3片，则上图中的流水任务运行起来的示意图如下，从运行图中可以看出，对于同一片数据，Stage1、Stage2、Stage3之间的处理具有依赖关系，需要串行处理；不同的数据切片，同一时间点，可以有多个任务在并行处理，由此达到任务并行、提升性能的目的。
AscendC分别针对Vector、Cube编程设计了不同的流水任务。开发者只需要完成基本任务的代码实现即可，底层的指令同步和并行调度由AscendC框架实现，开发者无需关注。
2.2矢量编程范式
矢量编程范式把算子的实现流程分为3个基本任务：CopyIn，Compute，CopyOut。CopyIn负责搬入操作，Compute负责矢量计算操作，CopyOut负责搬出操作。
我们只需要根据编程范式完成基本任务的代码实现就可以了，底层的指令同步和并行调度由AscendC框架来实现。
那AscendC是怎么完成不同任务之间的数据通信和同步的呢？这里AscendC提供了Queue队列管理的API，主要就是两个队列操作APIEnQue、DeQue以及内存的逻辑抽象。
矢量编程中使用到的逻辑位置定义如下：
·搬入数据的存放位置：VECIN；
·计算中间变量的位置：VECCALC；
·搬出数据的存放位置：VECOUT。
从前面可以看到，矢量编程主要分为CopyIn、Compute、CopyOut三个任务。CopyIn任务中将输入数据从Global内存搬运至Local内存后，需要使用EnQue将LocalTensor放入VECIN的Queue中；Compute任务等待VECIN的Queue中LocalTensor出队之后才可以完成矢量计算，计算完成后使用EnQue将计算结果LocalTensor放入到VECOUT的Queue中；CopyOut任务等待VECOUT的Queue中LocalTensor出队，再将其拷贝到Global内存。这样，Queue队列就完成了三个任务间的数据通信和同步。具体流程和流程图如下：
1.Stage1：CopyIn任务。
使用DataCopy接口将GlobalTensor数据拷贝到LocalTensor。
使用EnQue接口将LocalTensor放入VECIN的Queue中。
2.Stage2：Compute任务。
使用DeQue接口从VECIN中取出LocalTensor。
使用AscendC接口完成矢量计算。
使用EnQue接口将计算结果LocalTensor放入到VECOUT的Queue中。
3.Stage3：CopyOut任务。
使用DeQue接口从VECOUT的Queue中去除LocalTensor。
使用DataCopy接口将LocalTensor拷贝到GlobalTensor上。
这样我们的kernel实现代码就很清晰了。先初始化内存和队列，然后通过编程范式实现CopyIn、Compute、CopyOut三个Stage就可以了。
2.3SPMD并行编程-多核
最前面介绍昇腾AI处理器的时候，有介绍过AICore是有多个的，那我们怎么把多个AICore充分利用起来呢？常用的并行计算方法中，有一种SPMD数据并行的方法，简单说就是将数据分片，每片数据经过完整的一个数据处理流程。这个就能和昇腾AI处理器的多核匹配上了，我们将数据分成多份，每份数据的处理运行在一个核上，这样每份数据并行处理完成，整个数据也就处理完了。AscendC是SPMD编程，多个AICore共享相同的指令代码，每个核上的运行实例唯一的区别是就是block_idx不同，这样我们就可以通过block_idx来区分不同的核，只要对GlobalMemory上的数据地址进行切分偏移，就可以让每个核处理自己对应的那部分数据了。
算子被调用时，所有的计算核心都执行相同的实现代码，入口函数的入参也是相同的。每个核上处理的数据地址需要在起始地址上增加block_idx*BLOCK_LENGTH的偏移来获取。这样也就实现了多核并行计算的数据切分。
classKernelAdd{
public:
__aicore__inlineKernelAdd(){}
__aicore__inlinevoidInit(GM_ADDRx,GM_ADDRy,GM_ADDRz)
{
//getstartindexforcurrentcore,coreparallel
GM_ADDRxGmOffset=x+BLOCK_LENGTH*GetBlockIdx();
GM_ADDRyGmOffset=y+BLOCK_LENGTH*GetBlockIdx();
GM_ADDRzGmOffset=z+BLOCK_LENGTH*GetBlockIdx();
xGm.SetGlobalBuffer((__gm__half*)xGmOffset,BLOCK_LENGTH);
yGm.SetGlobalBuffer((__gm__half*)yGmOffset,BLOCK_LENGTH);
zGm.SetGlobalBuffer((__gm__half*)zGmOffset,BLOCK_LENGTH);
……
}
……
}
2.4AscendCAPI介绍
在整个kernel实现中，最最核心的代码就是Add(zLocal,xLocal,yLocal,TILE_LENGTH);通过一个AscendC提供的API接口完成了所有数据的加法计算，对，没看错，就是这个接口完成了计算。
接下来就介绍下AscendC提供的API。AscendC算子采用标准C++语法和一组类库API进行编程，类库API主要包含以下几种，大家可以在核函数的实现中根据自己的需求选择合适的API：
·计算类API，包括标量计算API、向量计算API、矩阵计算API，分别实现调用Scalar计算单元、Vector计算单元、Cube计算单元执行计算的功能。
·数据搬运API，上述计算API基于LocalMemory数据进行计算，所以数据需要先从GlobalMemory搬运至LocalMemory，再使用计算接口完成计算，最后从LocalMemory搬出至GlobalMemory。执行搬运过程的接口称之为数据搬移接口，比如DataCopy接口。
·内存管理API，用于分配管理内存，比如AllocTensor、FreeTensor接口。
·任务同步API，完成任务间的通信和同步，比如EnQue、DeQue接口。
AscendCAPI的计算操作数都是Tensor类型：GlobalTensor和LocalTensor。
介绍完AscendCAPI种类后，下面来解释下为什么一个Add接口就可以计算所有的数。原来AscendC编程模型是基于SIMD架构的，单条指令可以完成多个数据操作，同时在API内部封装了一些指令的高级功能。
2.5算子执行基本流程
前面有提到，在异构计算架构中，NPU与CPU是协同工作的，在AscendC编程模型中，我们需要实现NPU侧的代码和CPU侧的代码。在NPU侧的代码我们通常叫做Kernel实现代码，CPU侧的代码我们一般叫做Host实现代码，一份完整的AscendC代码，通常包括Host侧实现代码和Kernel侧实现代码。AscendC算子执行的基本流程如下：
1、初始化Device设备；
2、创建Context绑定设备；
3、分配Host内存，并进行数据初始化；
4、分配Device内存，并将数据从Host上拷贝到Device上；
5、用内核调用符调用核函数完成指定的运算；
6、将Device上的运算结果拷贝回Host；
7、释放申请的资源。
2.6核函数介绍
上面的流程中，最重要的一步就是调用核函数来进行并行计算任务。核函数是AscendC算子Device侧实现的入口。在核函数中，需要为在AI核上执行的代码规定要进行的数据访问和计算操作。
externC__global____aicore__voidadd_custom(__gm__uint8_t*x,__gm__uint8_t*y,__gm__uint8_t*z);
上面这个是一个核函数声明的示例，externC表示核函数按照类C的编译和连接规约来编译和连接，__global__函数类型限定符表示它是一个核函数，__aicore__函数类型限定符表示该核函数在device侧的AICore上执行。参数列表中的变量类型限定符__gm__，表明该指针变量指向GlobalMemory上某处内存地址，注意这里的入参只能支持指针或C/C++内置数据类型，样例里指针使用的类型为uint8_t，在后续的使用中需要将其转化为实际的指针类型。
AscendC编程模型中的核函数采用内核调用符...来调用，样例如下：
kernel_name(argumentlist);
kernel_name即为上面讲的核函数名称，argumentlist是核函数的函数入参，在中间，有3个参数：
·blockDim，规定了核函数将会在几个核上执行，我们可以先设置为1；
·l2ctrl，保留参数，暂时设置为固定值nullptr，我们不用关注；
·stream，使用aclrtCreateStream创建，用于多线程调度。
3样例开发讲解
3.1样例代码结构
|--CMakeLists.txt//编译工程文件
|--cmake//编译工程文件
|--data_utils.h//数据读入写出函数
|--input//存放脚本生成的输入数据目录
|--leakyrelu_custom.cpp//算子kernel实现
|--leakyrelu_custom.py//输入数据和真值数据生成脚本文件
|--leakyrelu_custom_tiling.h//host侧tiling函数
|--main.cpp//主函数，host侧调用代码，含cpu域及npu域调用
|--output//存放算子运行输出数据和标杆数据的目录
|--readme.md//执行命令说明
|--run.sh//运行脚本
3.2主要文件
3.2.1输入数据和真值数据生成脚本文件：KERNEL_NAME.py。
根据算子的输入输出编写生成输入数据和真值数据的脚本。
本例子生成8*200*1024大小的fp16数据：
……
defgen_golden_data_simple():
total_length_imm=8*200*1024
tile_num_imm=8
//生成tilling的bin文件
total_length=np.array(total_length_imm,dtype=np.uint32)
tile_num=np.array(tile_num_imm,dtype=np.uint32)
scalar=np.array(0.1,dtype=np.float32)
tiling=(total_length,tile_num,scalar)
tiling_data=b''.join(x.tobytes()forxintiling)
withos.fdopen(os.open('./input/tiling.bin',WRITE_FILE_FLAGS,PEN_FILE_MODES_640),'wb')asf:
f.write(tiling_data)
//生成输入数据
input_x=np.random.uniform(-100,100,[8,200,1024]).astype(np.float16)
//生成golden数据，功能和LeakyRelu相同
golden=np.where(input_x0,input_x,input_x*scalar).astype(np.float16)
input_x.tofile(./input/input_x.bin)
golden.tofile(./output/golden.bin)
3.2.2编译工程文件：CMakeLists.txt
用于编译cpu侧或npu侧运行的AscendC算子。主要关注CMakeLists.txt中源文件是否全部列全。
3.2.3调用算子的应用程序：main.cpp
主要是内存申请，数据拷贝和文件读写等操作，并最终调用算子，相关API的介绍如下：
1、AscendCL初始化接口aclInit，用于运行时接口AscendCL的初始化，是程序最先调用的接口；aclrtCreateContext和aclrtCreateStream用于创建Context和Stream，主要用于线程相关的资源管理。
2、aclrtMallocHost接口，用于在Host上申请内存：
aclErroraclrtMallocHost(void**hostPtr,size_tsize)
这个函数和C语言中的malloc类似，用于在Host上申请一定字节大小的内存，其中hostPtr是指向所分配内存的指针，size是申请的内存大小，如果需要释放这块内存的话，使用aclrtFreeHost接口释放，这和C语言中的free函数对应。
3、aclrtMalloc接口，用于在Device上申请内存：
aclErroraclrtMalloc(void**devPtr,size_tsize,aclrtMemMallocPolicypolicy)
和Host上的内存申请接口相比，多了一个policy参数，用于设置内存分配规则，一般设置成ACL_MEM_MALLOC_HUGE_FIRST就可以了。使用完毕后可以用对应的aclrtFree接口释放内存。
4、aclrtMemcpy接口，用于Host和Device之间数据拷贝：
前面申请的内存区分了Host内存和Device内存，那就会涉及到数据同步的问题，aclrtMemcpy就是用于Host和Device之间数据通信的接口：
aclErroraclrtMemcpy(void*dst,size_tdestMax,constvoid*src,size_tcount,aclrtMemcpyKindkind)
其中src指向数据源，而dst是目标内存地址，destMax是目的内存地址的最大内存长度，count是拷贝的字节数，其中aclrtMemcpyKind控制复制的方向：ACL_MEMCPY_HOST_TO_HOST、ACL_MEMCPY_HOST_TO_DEVICE、ACL_MEMCPY_DEVICE_TO_HOST和ACL_MEMCPY_DEVICE_TO_DEVICE，像ACL_MEMCPY_HOST_TO_DEVICE就是将Host上数据拷贝到Device上。
5、核心函数为CPU侧的调用kernel函数
ICPU_RUN_KF(leakyrelu_custom,blockDim,x,y,usrWorkSpace,tiling);
和NPU侧调用的
leakyrelu_custom_do(blockDim,nullptr,stream,xDevice,yDevice,workspaceDevice,tilingDevice);
完整代码如下：
//Thisfileconstainscodeofcpudebugandnpucode.Wereaddatafrombinfileandwriteresulttofile.
#includedata_utils.h
#includeleakyrelu_custom_tiling.h
#ifndef__CCE_KT_TEST__
#includeacl/acl.h
externvoidleakyrelu_custom_do(uint32_tcoreDim,void*l2ctrl,void*stream,uint8_t*x,uint8_t*y,
uint8_t*workspace,uint8_t*tiling);
#else
#includetikicpulib.h
externC__global____aicore__voidleakyrelu_custom(GM_ADDRx,GM_ADDRy,GM_ADDRworkspace,GM_ADDRtiling);
#endif
int32_tmain(int32_targc,char*argv[])
{
size_ttilingSize=sizeof(LeakyReluCustomTilingData);
size_tusrWorkspaceSize=4096;
size_tsysWorkspaceSize=16*1024*1024;
uint32_tblockDim=8;
#ifdef__CCE_KT_TEST__//CPU侧调用
//申请内存用于存放workspace和tilling数据
uint8_t*usrWorkSpace=(uint8_t*)AscendC::GmAlloc(usrWorkspaceSize);
uint8_t*tiling=(uint8_t*)AscendC::GmAlloc(tilingSize);
ReadFile(./input/tiling.bin,tilingSize,tiling,tilingSize);
size_tinputByteSize=blockDim*200*1024*sizeof(uint16_t);//uint16_trepresenthalf
size_toutputByteSize=blockDim*200*1024*sizeof(uint16_t);//uint16_trepresenthalf
//申请内存用于存放输入和输出数据
uint8_t*x=(uint8_t*)AscendC::GmAlloc(inputByteSize);
uint8_t*y=(uint8_t*)AscendC::GmAlloc(inputByteSize);
//获取输入数据
ReadFile(./input/input_x.bin,inputByteSize,x,inputByteSize);
//PrintData(x,16,printDataType::HALF);
//在AIV上执行
AscendC::SetKernelMode(KernelMode::AIV_MODE);
//调用kernel函数
ICPU_RUN_KF(leakyrelu_custom,blockDim,x,y,usrWorkSpace,tiling);//usethismacroforcpudebug
//PrintData(y,16,printDataType::HALF);
WriteFile(./output/output_y.bin,y,outputByteSize);
AscendC::GmFree((void*)x);
AscendC::GmFree((void*)y);
AscendC::GmFree((void*)usrWorkSpace);
AscendC::GmFree((void*)tiling);
#else//NPU侧调用
CHECK_ACL(aclInit(nullptr));
aclrtContextcontext;
int32_tdeviceId=0;
CHECK_ACL(aclrtSetDevice(deviceId));
CHECK_ACL(aclrtCreateContext(context,deviceId));
aclrtStreamstream=nullptr;
CHECK_ACL(aclrtCreateStream(stream));
uint8_t*xHost,*yHost,*tilingHost,*workspaceHost;
uint8_t*xDevice,*yDevice,*tilingDevice,*workspaceDevice;
//申请host上tilling内存并读入tilling数据
CHECK_ACL(aclrtMallocHost((void**)(tilingHost),tilingSize));
ReadFile(./input/tiling.bin,tilingSize,tilingHost,tilingSize);
//申请host上workspace内存
CHECK_ACL(aclrtMallocHost((void**)(workspaceHost),tilingSize));
size_tinputByteSize=blockDim*200*1024*sizeof(uint16_t);//uint16_trepresenthalf
size_toutputByteSize=blockDim*200*1024*sizeof(uint16_t);//uint16_trepresenthalf
size_tworkspaceByteSize=sysWorkspaceSize+usrWorkspaceSize;
//申请host和device上的输入输出内存和device上的workspace和tilling内存
CHECK_ACL(aclrtMallocHost((void**)(xHost),inputByteSize));
CHECK_ACL(aclrtMallocHost((void**)(yHost),inputByteSize));
CHECK_ACL(aclrtMallocHost((void**)(workspaceHost),workspaceByteSize));
CHECK_ACL(aclrtMalloc((void**)xDevice,inputByteSize,ACL_MEM_MALLOC_HUGE_FIRST));
CHECK_ACL(aclrtMalloc((void**)yDevice,inputByteSize,ACL_MEM_MALLOC_HUGE_FIRST));
CHECK_ACL(aclrtMalloc((void**)tilingDevice,tilingSize,ACL_MEM_MALLOC_HUGE_FIRST));
CHECK_ACL(aclrtMalloc((void**)workspaceDevice,workspaceByteSize,ACL_MEM_MALLOC_HUGE_FIRST));
ReadFile(./input/input_x.bin,inputByteSize,xHost,inputByteSize);
//PrintData(xHost,16,printDataType::HALF);
//从host上拷贝输入数据和tilling数据到device
CHECK_ACL(aclrtMemcpy(xDevice,inputByteSize,xHost,inputByteSize,ACL_MEMCPY_HOST_TO_DEVICE));
CHECK_ACL(aclrtMemcpy(tilingDevice,tilingSize,tilingHost,tilingSize,ACL_MEMCPY_HOST_TO_DEVICE));
//调用核函数
leakyrelu_custom_do(blockDim,nullptr,stream,xDevice,yDevice,workspaceDevice,tilingDevice);
//等待核函数运行完成
CHECK_ACL(aclrtSynchronizeStream(stream));
//拷回运行结果到host
CHECK_ACL(aclrtMemcpy(yHost,outputByteSize,yDevice,outputByteSize,ACL_MEMCPY_DEVICE_TO_HOST));
//PrintData(yHost,16,printDataType::HALF);
WriteFile(./output/output_y.bin,yHost,outputByteSize);
//释放资源
CHECK_ACL(aclrtFree(xDevice));
CHECK_ACL(aclrtFree(yDevice));
CHECK_ACL(aclrtFree(workspaceDevice));
CHECK_ACL(aclrtFree(tilingDevice));
CHECK_ACL(aclrtFreeHost(xHost));
CHECK_ACL(aclrtFreeHost(yHost));
CHECK_ACL(aclrtFreeHost(workspaceHost));
CHECK_ACL(aclrtFreeHost(tilingHost));
CHECK_ACL(aclrtDestroyStream(stream));
CHECK_ACL(aclrtDestroyContext(context));
CHECK_ACL(aclrtResetDevice(deviceId));
CHECK_ACL(aclFinalize());
#endif
return0;
}
3.2.4一键式编译运行脚本run.sh
编译和运行应用程序。
cpu侧运行命令：
bashrun.shleakyrelu_customascend910B1VectorCorecpu
npu侧运行命令：
bashrun.shleakyrelu_customascend910B1VectorCorenpu
参数含义如下：
bashrun.sh
表示需要运行的算子。表示算子运行的AI处理器型号。表示在AICore上或者VectorCore上运行，参数取值为AiCore/VectorCore。表示算子以cpu模式或npu模式运行，参数取值为cpu/npu。
3.3kernel实现
3.3.1函数原型定义
本样例中，函数名为leakyrelu_custom，根据对算子输入输出的分析，确定有2个参数x，y，其中x为输入内存，y为输出内存。核函数原型定义如下所示：
externC__global____aicore__voidleakyrelu_custom(GM_ADDRx,GM_ADDRy,GM_ADDRworkspace,GM_ADDRtiling){}
使用__global__函数类型限定符来标识它是一个核函数，可以被...调用；使用__aicore__函数类型限定符来标识该核函数在设备端AICore上执行；为方便起见，统一使用GM_ADDR宏修饰入参，GM_ADDR宏定义：
#defineGM_ADDR__gm__uint8_t*__restrict__
3.3.2获取tilling数据，并调用算子类的Init和Process函数。
算子类的Init函数，完成内存初始化相关工作，Process函数完成算子实现的核心逻辑。
externC__global____aicore__voidleakyrelu_custom(GM_ADDRx,GM_ADDRy,GM_ADDRworkspace,GM_ADDRtiling)
{
GET_TILING_DATA(tilingData,tiling);
KernelLeakyReluop;
op.Init(x,y,tilingData.totalLength,tilingData.tileNum,tilingData.scalar);
op.Process();
}
3.3.3对核函数的调用进行封装
封装后得到leakyrelu_custom_do函数，便于主程序调用。#ifndef__CCE_KT_TEST__表示该封装函数仅在编译运行NPU侧的算子时会用到，编译运行CPU侧的算子时，可以直接调用add_custom函数。调用核函数时，除了需要传入输入输出参数x，y，切分相关参数tiling，还需要传入blockDim,l2ctrl,stream来规定核函数的执行配置。
#ifndef__CCE_KT_TEST__
//callofkernelfunction
voidleakyrelu_custom_do(uint32_tblockDim,void*l2ctrl,void*stream,uint8_t*x,uint8_t*y,
uint8_t*workspace,uint8_t*tiling)
{
leakyrelu_custom(x,y,workspace,tiling);
}
#endif
3.3.4获取tiling参数
主要从tilingPointer中获取tiling的参数totalLength(总长度)、tileNum和scalar。
#defineGET_TILING_DATA(tilingData,tilingPointer)
LeakyReluCustomTilingDatatilingData;
INIT_TILING_DATA(LeakyReluCustomTilingData,tilingDataPointer,tilingPointer);
(tilingData).totalLength=tilingDataPointer-totalLength;
(tilingData).tileNum=tilingDataPointer-tileNum;
(tilingData).scalar=tilingDataPointer-scalar;
#endif//LEAKYRELU_CUSTOM_TILING_H
3.3.5Init函数
主要获取tiling数据后，设置单核上gm的地址和Buffer的初始化。
__aicore__inlinevoidInit(GM_ADDRx,GM_ADDRy,uint32_ttotalLength,uint32_ttileNum,floatscalar)
{
ASSERT(GetBlockNum()!=0blockdimcannotbezero!);
this-blockLength=totalLength/GetBlockNum();
this-tileNum=tileNum;
this-scalar=static_cast(scalar);
ASSERT(tileNum!=0tilenumcannotbezero!);
this-tileLength=this-blockLength/tileNum/BUFFER_NUM;
//getstartindexforcurrentcore,coreparallel
xGm.SetGlobalBuffer((__gm__half*)x+this-blockLength*get_block_idx(),this-blockLength);
yGm.SetGlobalBuffer((__gm__half*)y+this-blockLength*get_block_idx(),this-blockLength);
//pipeallocmemorytoqueue,theunitisBytes
pipe.InitBuffer(inQueueX,BUFFER_NUM,this-tileLength*sizeof(half));
pipe.InitBuffer(outQueueY,BUFFER_NUM,this-tileLength*sizeof(half));
}
3.3.6Process函数
主要实现三个CopyIn、Compute、CopyOut这三stage。
__aicore__inlinevoidProcess()
{
//loopcountneedtobedoubled,duetodoublebuffer
int32_tloopCount=this-tileNum*BUFFER_NUM;
//tilingstrategy,pipelineparallel
for(int32_ti=0;iloopCount;i++){
CopyIn(i);
Compute(i);
CopyOut(i);
}
}
3.3.7CopyIn函数
负责从GlobalMemory拷贝数据到LocalMemory，并将数据加入Queue
__aicore__inlinevoidCopyIn(int32_tprogress)
{
//alloctensorfromqueuememory
LocalTensorxLocal=inQueueX.AllocTensor();
//copyprogress_thtilefromglobaltensortolocaltensor
DataCopy(xLocal,xGm[progress*tileLength],tileLength);
//enqueinputtensorstoVECINqueue
inQueueX.EnQue(xLocal);
}
3.3.8Compute函数
负责从Queue中取出数据，进行计算，并将结果放入Queue
__aicore__inlinevoidCompute(int32_tprogress)
{
//dequeinputtensorsfromVECINqueue
LocalTensorxLocal=inQueueX.DeQue();
LocalTensoryLocal=outQueueY.AllocTensor();
//callLeakyReluinstrforcomputation
LeakyRelu(yLocal,xLocal,scalar,tileLength);
//enquetheoutputtensortoVECOUTqueue
outQueueY.EnQue(yLocal);
//freeinputtensorsforreuse
inQueueX.FreeTensor(xLocal);
}
3.3.9CopyOut函数
负责从Queue中将数据取出，并将数据从LocalMemory拷贝到GlobalMemory。
__aicore__inlinevoidCopyOut(int32_tprogress)
{
//dequeoutputtensorfromVECOUTqueue
LocalTensoryLocal=outQueueY.DeQue();
//copyprogress_thtilefromlocaltensortoglobaltensor
DataCopy(yGm[progress*tileLength],yLocal,tileLength);
//freeoutputtensorforreuse
outQueueY.FreeTensor(yLocal);
}
3.4编译和执行
3.4.1在CPU侧执行
执行结果如下：
可以看到最后的输出结果output_y.bin和标杆数据golden.bin的MD5值相同，说明计算结果相同。
执行完成后，在input下存放输入数据和tiling数据，在output下面存放了输出数据和标杆数据，npuchk目录下是每个核的npu_check执行结果
在当前目录还有一个可执行二进制文件leakyrelu_custom_cpu，如果执行报错，可以通过gdb调试这个可执行文件，具体调试可参考文末官方教程。
3.4.2在NPU侧执行
在NPU侧执行有两种方式：仿真执行和上板运行，命令都相同，只是编译选项不同，我们可以通过修改编译选项-DASCEND_RUN_MODE为SIMULATOR运行CAModel仿真，设置为ONBOARD是上板运行。
functioncompile_and_execute(){
#使用cmake编译cpu侧或者npu侧算子,SIMULATORorONBOARD
mkdir-pbuild;cdbuild;
cmake..
-Dsmoke_testcase=$1
-DASCEND_PRODUCT_TYPE=$2
-DASCEND_CORE_TYPE=$3
-DASCEND_RUN_MODE=SIMULATOR
-DASCEND_INSTALL_PATH=$ASCEND_HOME_DIR
VERBOSE=1cmake--build.--target${1}_${4}
……
}
4参考资料
总之，学习AscendC，仅需了解C++编程、理解对列通信与内存申请释放机制、通过调用相应的计算接口与搬运接口，就可以写出运行在昇腾AI处理器上的高性能算子。
了解更多AscendC学习资源，请访问官方教程：AscendC编程指南